security: harden model-file parsing, server, and build (untrusted-mirror threat model)

c/Makefile

@@ -148,16 +148,32 @@ else
 CUDA_HOME ?= /usr/local/cuda
 NVCC      ?= $(CUDA_HOME)/bin/nvcc
 endif
+# CUDA target selection. Default `native` builds SASS for the build machine's GPU
+# only (fast local dev). A headless CI/Docker build has no GPU, so `-arch=native`
+# there detects nothing and the shipped coli_cuda.dll fails at first kernel dispatch.
+# For a portable/release DLL that runs on any Ampere..Blackwell card AND JITs on
+# newer archs, build with CUDA_ARCH=portable: it emits SASS for sm_80/86/89/90/120
+# plus a compute_120 PTX fallback (sm_120 / RTX 50-series needs CUDA >= 12.8).
 CUDA_ARCH ?= native
+ifeq ($(CUDA_ARCH),portable)
+CUDA_GENCODE = -gencode arch=compute_80,code=sm_80 \
+               -gencode arch=compute_86,code=sm_86 \
+               -gencode arch=compute_89,code=sm_89 \
+               -gencode arch=compute_90,code=sm_90 \
+               -gencode arch=compute_120,code=sm_120 \
+               -gencode arch=compute_120,code=compute_120
+else
+CUDA_GENCODE = -arch=$(CUDA_ARCH)
+endif
 ifneq ($(IS_WIN),)
 # nvcc's host compiler on Windows is MSVC cl.exe: the GCC-style
 # -Xcompiler=-Wall,-Wextra is rejected ("D8021 invalid numeric argument
 # '/Wextra'"), which made `make cuda-dll` unbuildable as shipped. -W3 is
 # the MSVC warning level (dash form, NOT /W3: MSYS make would mangle the
 # slash-form into a filesystem path).
-NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-W3
+NVCCFLAGS ?= -O3 -std=c++17 $(CUDA_GENCODE) -Xcompiler=-W3
 else
-NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-Wall,-Wextra
+NVCCFLAGS ?= -O3 -std=c++17 $(CUDA_GENCODE) -Xcompiler=-Wall,-Wextra
 endif
 # HIP=1 builds the SAME backend for AMD GPUs via ROCm: backend_cuda.cu is
 # compiled unchanged through backend_gpu_compat.h (one source, two vendors,

c/backend_cuda.cu

@@ -793,7 +793,8 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
     if(!cuda_ok(copy_desc,"expert group descriptors"))return 0;
     int profile=getenv("COLI_CUDA_PROFILE")&&atoi(getenv("COLI_CUDA_PROFILE"));
     cudaEvent_t ev[4]={};
-    if(profile) for(int i=0;i<4;i++) if(!cuda_ok(cudaEventCreate(&ev[i]),"profile event")) profile=0;
+    if(profile) for(int i=0;i<4;i++) if(!cuda_ok(cudaEventCreate(&ev[i]),"profile event")){
+        for(int j=0;j<i;j++) cudaEventDestroy(ev[j]); profile=0; break; }   /* (#B8) don't leak the events already created */
     if(profile) cudaEventRecord(ev[0],ctx->stream);
     if(async)std::memcpy(ctx->host_x,x,xb);
     cudaError_t copy_x=async?cudaMemcpyAsync(ctx->x,ctx->host_x,xb,cudaMemcpyHostToDevice,ctx->stream)

c/colibri.c

@@ -819,6 +819,21 @@ static void rope_interleave(float *v, int pos, const Cfg *c){
 }
 
 /* ---------- config ---------- */
+/* SEC-9: bounded slurp for untrusted config/oracle JSON. config.json arrives from
+ * unverified mirrors (see qt_check_fmt threat model); an unbounded ftell->malloc
+ * gave a hostile file a load-time OOM or, on malloc failure, a NULL deref via
+ * b[got]=0. Cap the size, NULL-check the alloc, require a full read. Returns a
+ * malloc'd NUL-terminated buffer, or NULL on any failure. Mirrors tok.h tk_read_file. */
+#define CFG_MAX_BYTES (256ll<<20)   /* config/oracle JSON is KB-MB in practice */
+static char* cfg_slurp(const char *path){
+    FILE *f=fopen(path,"rb"); if(!f) return NULL;
+    fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET);
+    if(n<0 || (long long)n>CFG_MAX_BYTES){ fclose(f); return NULL; }
+    char *b=malloc((size_t)n+1); if(!b){ fclose(f); return NULL; }
+    size_t got=fread(b,1,(size_t)n,f); fclose(f);
+    if((long)got!=n){ free(b); return NULL; }
+    b[got]=0; return b;
+}
 static jval* cfg_root(const char *snap, char **arena){
     char p[2048]; snprintf(p,sizeof(p),"%s/config.json",snap);
     FILE *f=fopen(p,"rb"); if(!f){perror(p);exit(1);}
@@ -987,11 +1002,18 @@ static void qt_from_disk(Model *m, const char *name, int O, int I, int bits, int
             if(t->fmt!=5||!t->q4){ t->fmt=5; t->O=O; t->I=I; t->gs=0; t->q4=qalloc(nb); t->s=falloc((int64_t)O*ng); }
             st_read_raw(&m->S,name,t->q4,drop); }
         else      { if(t->fmt!=fmt||!t->q4){ t->fmt=fmt; t->O=O; t->I=I; t->gs=0; t->q4=qalloc(nb); t->s=qsalloc(O); } st_read_raw(&m->S,name,t->q4,drop); }
-        st_read_f32(&m->S,sn,t->s,drop);
+        /* cap MUST match the scale cardinality qt_resolve_fmt already validated and
+         * the falloc above actually reserved, per format: grouped-int4 (fmt=4) keeps
+         * O*ceil(I/gs) scales and int3-g64 (fmt=5) keeps O*i3_groups(I); everything
+         * else is per-row O. Using the per-row bound for a grouped format would
+         * reject a legitimate container (fmt=5 regressed exactly that way). */
+        st_read_f32_cap(&m->S,sn,t->s,
+                        fmt==4 ? (int64_t)O*((I+gs-1)/gs) :
+                        fmt==5 ? (int64_t)O*i3_groups(I)  : (int64_t)O, drop);
     } else {
         if(!t->qf && !t->q8 && !t->q4) qt_alloc(t,O,I,bits);
-        if(t->fmt==0) st_read_f32(&m->S,name,t->qf,drop);
-        else { float *tmp=falloc((int64_t)O*I); st_read_f32(&m->S,name,tmp,drop); qt_fill(t,tmp,bits); free(tmp); }
+        if(t->fmt==0) st_read_f32_cap(&m->S,name,t->qf,(int64_t)O*I,drop);
+        else { float *tmp=falloc((int64_t)O*I); st_read_f32_cap(&m->S,name,tmp,(int64_t)O*I,drop); qt_fill(t,tmp,bits); free(tmp); }
     }
 }
 static QT qt_load(Model *m, const char *name, int O, int I, int bits){
@@ -1213,6 +1235,7 @@ static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits
 /* embed: dequantizza la riga del token (scala per-riga) in x[hidden] */
 static void embed_row(Model *m, int tok, float *x){
     int D=m->c.hidden; QT *e=&m->embed;
+    if(tok<0 || tok>=e->O){ memset(x,0,(size_t)D*sizeof(float)); return; }   /* #SEC-5: out-of-range token id -> zero row, never OOB */
     if(e->fmt==0){ memcpy(x, e->qf+(int64_t)tok*D, D*sizeof(float)); return; }
     if(e->fmt==4){ /* grouped int4: per-group scale (embed/lm_head at io_bits, usually fmt 0/1) */
         const uint8_t *q=e->q4+(int64_t)tok*((D+1)/2); int gs=e->gs,ng=(D+gs-1)/gs;
@@ -6549,10 +6572,8 @@ int main(int argc, char **argv){
 
     /* altrimenti: validazione contro l'oracolo (ref_glm.json) */
     const char *refpath=getenv("REF")?getenv("REF"):"ref_glm.json";
-    FILE *f=fopen(refpath,"rb"); if(!f){perror(refpath);return 1;}
-    fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET);
-    char *b=malloc(n+1); size_t got=fread(b,1,n,f); b[got]=0; fclose(f);
-    if((long)got!=n) fprintf(stderr,"warning: short read on %s (%ld of %ld)\n",refpath,(long)got,n);
+    char *b=cfg_slurp(refpath);
+    if(!b){ fprintf(stderr,"%s: cannot read oracle file (missing, unreadable, short, or > %lld bytes)\n",refpath,(long long)CFG_MAX_BYTES); return 1; }
     char *ar=NULL; jval *ref=json_parse(b,&ar);
     int np=0,nfull=0; int *prompt=read_arr(ref,"prompt_ids",&np); int *full=read_arr(ref,"full_ids",&nfull);
     if(!prompt||!full||np<1||nfull<np){ fprintf(stderr,"ref file missing prompt_ids/full_ids or empty\n"); return 1; }

c/download_fp8.py

@@ -53,14 +53,14 @@ def download_file_ms(fn):
         model_id=REPO_MS,
         file_path=fn,
         local_dir=DEST,
-        revision="master",
+        revision=os.environ.get("GLM_MS_REVISION", "master"),   # pin to a commit for supply-chain integrity
     )
 
 def download_file_hf(fn):
     """Download a single file from HuggingFace with hf_transfer."""
     os.environ["HF_HUB_ENABLE_HF_TRANSFER"] = "1"
     from huggingface_hub import hf_hub_download
-    hf_hub_download(REPO_HF, fn, local_dir=DEST)
+    hf_hub_download(REPO_HF, fn, local_dir=DEST, revision=os.environ.get("GLM_HF_REVISION", "main"))
 
 def download_file_curl(fn, base_url, expected_size):
     """Fallback: download with curl to a .part file with resume."""

c/json.h

@@ -55,10 +55,12 @@ static jval *j_parse_val(jparser *p);
 static char *j_parse_str_raw(jparser *p) {
     /* assume *p->s == '"' */
     p->s++;
-    const char *start = p->s;
-    /* trova la fine gestendo gli escape, poi copia decodificando i casi base */
-    char tmp[1 << 16]; int n = 0;
-    #define J_PUT(ch) do{ if (n < (int)sizeof(tmp)-1) tmp[n++] = (char)(ch); }while(0)
+    /* buffer su heap che CRESCE: niente troncamento silenzioso a 64KB (le stringhe
+     * lunghe di tokenizer.json/config venivano tagliate) e niente 64KB di stack. */
+    size_t cap = 64, n = 0; char *tmp = (char *)malloc(cap);
+    if (!tmp) { fprintf(stderr, "OOM parsing JSON string\n"); exit(1); }
+    #define J_PUT(ch) do{ if (n + 1 >= cap) { cap *= 2; tmp = (char *)realloc(tmp, cap); \
+        if (!tmp) { fprintf(stderr, "OOM parsing JSON string\n"); exit(1); } } tmp[n++] = (char)(ch); }while(0)
     while (*p->s && *p->s != '"') {
         char c = *p->s++;
         if (c == '\\' && *p->s) {
@@ -69,9 +71,11 @@ static char *j_parse_str_raw(jparser *p) {
                 case 'f': c = '\f'; break; case '/': c = '/'; break;
                 case '\\': c = '\\'; break; case '"': c = '"'; break;
                 case 'u': {  /* \uXXXX -> codepoint UTF-8 (con coppie surrogate) */
+                    if (!p->s[0]||!p->s[1]||!p->s[2]||!p->s[3]) { c='?'; break; }   /* \u troncato: non leggere oltre il NUL */
                     unsigned cp = (unsigned)strtoul((char[]){p->s[0],p->s[1],p->s[2],p->s[3],0}, NULL, 16);
                     p->s += 4;
-                    if (cp >= 0xD800 && cp <= 0xDBFF && p->s[0]=='\\' && p->s[1]=='u') {
+                    if (cp >= 0xD800 && cp <= 0xDBFF && p->s[0]=='\\' && p->s[1]=='u'
+                        && p->s[2] && p->s[3] && p->s[4] && p->s[5]) {
                         unsigned lo = (unsigned)strtoul((char[]){p->s[2],p->s[3],p->s[4],p->s[5],0}, NULL, 16);
                         if (lo >= 0xDC00 && lo <= 0xDFFF) { cp = 0x10000 + ((cp-0xD800)<<10) + (lo-0xDC00); p->s += 6; }
                     }
@@ -88,8 +92,8 @@ static char *j_parse_str_raw(jparser *p) {
     }
     #undef J_PUT
     if (*p->s == '"') p->s++;
-    (void)start;
-    return j_dup(p, tmp, n);
+    char *out = j_dup(p, tmp, (int)n); free(tmp);
+    return out;
 }
 
 static jval *j_parse_val(jparser *p) {
@@ -135,9 +139,9 @@ static jval *j_parse_val(jparser *p) {
         p->depth--;
         return v;
     }
-    if (c == 't') { p->s += 4; jval *v = j_new(J_BOOL); v->boolean = 1; return v; }
-    if (c == 'f') { p->s += 5; jval *v = j_new(J_BOOL); v->boolean = 0; return v; }
-    if (c == 'n') { p->s += 4; return j_new(J_NULL); }
+    if (c == 't' && !strncmp(p->s, "true", 4))  { p->s += 4; jval *v = j_new(J_BOOL); v->boolean = 1; return v; }
+    if (c == 'f' && !strncmp(p->s, "false", 5)) { p->s += 5; jval *v = j_new(J_BOOL); v->boolean = 0; return v; }
+    if (c == 'n' && !strncmp(p->s, "null", 4))  { p->s += 4; return j_new(J_NULL); }
     /* numero */
     { char *end; double d = strtod(p->s, &end); p->s = end; jval *v = j_new(J_NUM); v->num = d; return v; }
 }

c/olmoe.c

@@ -216,20 +216,35 @@ static void softmax_row(float *x, int n) {
 }
 
 /* ---------- caricamento ---------- */
+/* config.json arrives from an untrusted mirror: a missing key was a NULL-deref
+ * (json_get(...)->num), so require each dimension present + numeric. */
+static double req_num(jval *r, const char *k){
+    jval *v=json_get(r,k);
+    if(!v||v->t!=J_NUM){ fprintf(stderr,"config.json: missing or non-numeric \"%s\"\n",k); exit(1); }
+    return v->num;
+}
 static void load_cfg(Cfg *c, const char *snap) {
     char path[2048]; snprintf(path, sizeof(path), "%s/config.json", snap);
     FILE *f = fopen(path, "rb"); if(!f){perror(path);exit(1);}
     fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET);
-    char *buf = malloc(n+1); if(fread(buf,1,n,f)!=(size_t)n){} buf[n]=0; fclose(f);
+    if(n<0 || n>(256L<<20)){ fprintf(stderr,"%s: config.json missing or larger than 256 MB\n",path); exit(1); }  /* SEC-9 */
+    char *buf = malloc((size_t)n+1); if(!buf){ fprintf(stderr,"OOM reading %s\n",path); exit(1); }
+    if(fread(buf,1,(size_t)n,f)!=(size_t)n){ fprintf(stderr,"%s: short read\n",path); exit(1); } buf[n]=0; fclose(f);
     char *arena=NULL; jval *r = json_parse(buf, &arena);
-    c->hidden    = (int)json_get(r,"hidden_size")->num;
-    c->n_layers  = (int)json_get(r,"num_hidden_layers")->num;
-    c->n_heads   = (int)json_get(r,"num_attention_heads")->num;
-    c->n_kv_heads= (int)json_get(r,"num_key_value_heads")->num;
-    c->n_experts = (int)json_get(r,"num_experts")->num;
-    c->topk      = (int)json_get(r,"num_experts_per_tok")->num;
-    c->inter     = (int)json_get(r,"intermediate_size")->num;
-    c->vocab     = (int)json_get(r,"vocab_size")->num;
+    c->hidden    = (int)req_num(r,"hidden_size");
+    c->n_layers  = (int)req_num(r,"num_hidden_layers");
+    c->n_heads   = (int)req_num(r,"num_attention_heads");
+    c->n_kv_heads= (int)req_num(r,"num_key_value_heads");
+    c->n_experts = (int)req_num(r,"num_experts");
+    c->topk      = (int)req_num(r,"num_experts_per_tok");
+    c->inter     = (int)req_num(r,"intermediate_size");
+    c->vocab     = (int)req_num(r,"vocab_size");
+    /* range-check so bad dims can't drive a later malloc(inter*hidden) / div-by-zero */
+    if(c->hidden<1||c->hidden>(1<<20) || c->n_heads<1||c->n_heads>(1<<16) ||
+       c->inter<1||c->inter>(1<<24) || c->vocab<1||c->vocab>(1<<24) ||
+       c->n_layers<1||c->n_layers>4096 || c->n_experts<1||c->n_experts>(1<<20) ||
+       c->n_kv_heads<1 || c->topk<1||c->topk>c->n_experts){
+        fprintf(stderr,"config.json: dimension out of range\n"); exit(1); }
     c->head_dim  = c->hidden / c->n_heads;
     jval *th = json_get(r,"rope_theta");  c->theta = th ? (float)th->num : 10000.f;
     jval *ep = json_get(r,"rms_norm_eps"); c->eps   = ep ? (float)ep->num : 1e-5f;