From 9efa266a6c2581690fc2a16df00d5ac307cea8b5 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Wed, 22 Jul 2026 12:06:23 +0200 Subject: [PATCH 01/22] convert: MiniMax-M3 checkpoint support (--arch m3, auto-detected) Maps the minimax_m3_vl VL checkpoint onto the GLM container scheme so the engine's MoE loader stays single-naming: strips the language_model. prefix, block_sparse_moe -> mlp (router gate.weight + e_score_correction_bias line up with the existing loader), Mixtral w1/w3/w2 -> gate/up/down_proj, drops the vision tower + the MSA index branch (a later pass, like DSA), keeps q/k/v_proj under the attn bits class, q_norm/k_norm as f32. Writes a flattened text_config as the container config.json and carries chat_template.jinja. Validated on a synthetic tiny checkpoint: full name flow + bit-exact int4-g64 vs quant_int4_grouped. --- c/tools/convert_fp8_to_int4.py | 116 ++++++++++++++++++++++++++++----- 1 file changed, 99 insertions(+), 17 deletions(-) diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 6b9c9a1d0..439af24c6 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -228,12 +228,49 @@ def classify(name, n_layers, keep_mtp=False, keep_idx=False): if name.endswith("o_proj.weight"): return "o" if name.endswith("kv_b_proj.weight"): return "kvb" if any(name.endswith(k) for k in ("q_a_proj.weight", "q_b_proj.weight", - "kv_a_proj_with_mqa.weight")): return "attn" + "kv_a_proj_with_mqa.weight", + # GQA family (MiniMax-M3): plain q/k/v projections + "self_attn.q_proj.weight", "self_attn.k_proj.weight", + "self_attn.v_proj.weight")): return "attn" if any(name.endswith(k) for k in ("mlp.gate_proj.weight", "mlp.up_proj.weight", "mlp.down_proj.weight")): return "dmlp" if name.endswith(".weight"): return "q" # fallback: other resident weights return "f32" +# ---------- MiniMax-M3 (minimax_m3_vl): name normalization to the GLM container scheme ---- +# The VL checkpoint nests the text model under `language_model.` and uses Mixtral-style +# expert names. Normalizing here keeps the C engine's MoE loader single-naming: +# language_model.model.X -> model.X (language_model.lm_head -> lm_head) +# .block_sparse_moe. -> .mlp. (router gate.weight + bias line up) +# .experts.E.{w1,w3,w2}.weight -> .experts.E.{gate_proj,up_proj,down_proj}.weight +# vision_tower / multi_modal_projector / patch_merge_mlp -> dropped (text-only port) +# .self_attn.index_* -> dropped (MSA index branch; separate pass later, +# like the GLM DSA indexer) +# Returns the container name, or None to skip the tensor entirely. +def m3_name(name): + if name.startswith(("vision_tower.", "multi_modal_projector.", "patch_merge_mlp.")): + return None + if ".self_attn.index_" in name: return None + if name.startswith("language_model."): name = name[len("language_model."):] + name = name.replace(".block_sparse_moe.", ".mlp.") + name = name.replace(".w1.weight", ".gate_proj.weight") + name = name.replace(".w3.weight", ".up_proj.weight") + name = name.replace(".w2.weight", ".down_proj.weight") + return name + +# Container config.json for M3: the engine reads a flat config, so hoist text_config +# to the top level (keeping generation ids reachable) and drop the vision blocks. +def write_m3_config(src_path, outdir): + cfg = json.loads(open(src_path).read()) + tc = cfg.get("text_config", {}) + flat = dict(tc) + flat["model_type"] = "minimax_m3" + flat["architectures"] = ["MiniMaxM3SparseForCausalLM"] + for k in ("torch_dtype", "transformers_version"): + if k in cfg: flat[k] = cfg[k] + with open(os.path.join(outdir, "config.json"), "w") as f: + json.dump(flat, f, indent=1) + # ---------- dequant NVFP4 (modelopt) di UN tensore expert -> f32 [O,I] ---------- def dequant_nvfp4(f, name): """NVFP4 di NVIDIA modelopt (quant_algo=NVFP4, quant_method=modelopt). @@ -330,15 +367,18 @@ def _e8_job(item): return name, q, s def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, - keep_mtp=False, keep_idx=False, group_size=0, bits_map=None): + keep_mtp=False, keep_idx=False, group_size=0, bits_map=None, arch="glm"): from safetensors import safe_open e8_jobs = [] # deferred: encoded in a pool after the scan with safe_open(path, framework="pt") as f: keys = set(f.keys()) for name in f.keys(): - kind = classify(name, n_layers, keep_mtp, keep_idx) + cname = m3_name(name) if arch == "m3" else name + if cname is None: continue + kind = classify(cname, n_layers, keep_mtp, keep_idx) if kind in ("skip", "consumed"): continue w = dequant(f, name, keys) + name = cname if kind == "f32": out_dict[name] = w.astype(np.float32) else: @@ -401,11 +441,11 @@ def _init_worker(proj_bits): def _convert_one(args): # Convert one shard in a worker; the main process writes the result, so shard numbering # and the atomic manifest stay serial and identical to --workers 1. - i, sp, n_layers, ebits, io_bits, xbits, keep_mtp, keep_idx, group_size, bits_map = args + i, sp, n_layers, ebits, io_bits, xbits, keep_mtp, keep_idx, group_size, bits_map, arch = args out = {} convert_shard(sp, out, n_layers, ebits, io_bits, xbits, keep_mtp=keep_mtp, keep_idx=keep_idx, - group_size=group_size, bits_map=bits_map) + group_size=group_size, bits_map=bits_map, arch=arch) return i, out def check_or_record_params(outdir, prefix, params): @@ -483,7 +523,13 @@ def main(): help="bits for gate_proj in routed experts. Default=xbits") ap.add_argument("--down-bits", type=_bits, default=None, help="bits for down_proj in routed experts. Default=xbits") - ap.add_argument("--n-layers", type=int, default=78) + ap.add_argument("--n-layers", type=int, default=None, # default: from --arch (glm 78, m3 60) + help="main-model layer count (default: 78 for glm, 60 for m3)") + ap.add_argument("--arch", choices=["auto", "glm", "m3"], default="auto", + help="checkpoint family: glm (GLM/DeepSeek MLA names, the default when no " + "language_model. prefix is seen) or m3 (MiniMax-M3 VL: strips the " + "language_model. prefix, maps block_sparse_moe/w1/w2/w3 onto the GLM " + "container scheme, drops the vision tower + MSA index branch)") ap.add_argument("--min-free-gb", type=float, default=20.0) ap.add_argument("--workers", type=int, default=1, help="Parallel worker processes for the local --indir conversion " @@ -548,6 +594,33 @@ def main(): if bits_map: print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items()))) + # --arch auto: detect the checkpoint family BEFORE the plan print. config.json + # model_type is authoritative; a shard-key peek ("language_model." prefix) is the + # fallback when no config is present (e.g. a bare shard dir). + if a.arch == "auto": + mt = None + try: + if a.indir and os.path.exists(os.path.join(a.indir, "config.json")): + mt = json.loads(open(os.path.join(a.indir, "config.json")).read()).get("model_type") + elif a.repo: + from huggingface_hub import hf_hub_download + mt = json.loads(open(hf_hub_download(a.repo, "config.json")).read()).get("model_type") + except Exception: + mt = None + if mt is None and a.indir: + sh = sorted(glob.glob(os.path.join(a.indir, "*.safetensors"))) + if sh: + from safetensors import safe_open + with safe_open(sh[0], framework="pt") as f: + if any(k.startswith("language_model.") for k in f.keys()): mt = "minimax_m3_vl" + a.arch = "m3" if (mt or "").startswith("minimax") else "glm" + print(f"[ARCH] auto-detected: {a.arch}" + (f" (model_type {mt})" if mt else "")) + if a.n_layers is None: + a.n_layers = 60 if a.arch == "m3" else 78 + if a.arch == "m3" and (a.mtp or a.indexer): + raise SystemExit("--arch m3: the checkpoint ships no MTP tensors, and the MSA index " + "branch pass is not implemented yet (dropped in the main pass, like DSA)") + # 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 @@ -555,7 +628,7 @@ def main(): 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: {source_label(a)} | " + print(f"[PLAN] mode: {mode} | arch: {a.arch} | source: {source_label(a)} | " f"experts {a.ebits}-bit, embed/lm_head {a.io_bits}-bit, x {a.xbits}-bit | {grp}") if a.selftest_nvfp4: @@ -679,7 +752,7 @@ def get_slice(s, n): return None # 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, - "proj_bits": dict(PROJ_BITS)} + "proj_bits": dict(PROJ_BITS), "arch": a.arch} prog_path = os.path.join(a.outdir, f".{prefix}progress.json") prog = {} if os.path.exists(prog_path): @@ -702,7 +775,7 @@ def get_slice(s, n): return None if a.workers and a.workers > 1: from multiprocessing import Pool _pending = [(i, sp, a.n_layers, a.ebits, a.io_bits, a.xbits, - a.mtp, a.indexer, a.group_size, bits_map) + a.mtp, a.indexer, a.group_size, bits_map, a.arch) for i, sp in enumerate(shards) if not _shard_already_done(done, os.path.basename(sp), a.outdir)] _pool = Pool(a.workers, initializer=_init_worker, initargs=(dict(PROJ_BITS),)) @@ -728,7 +801,7 @@ def get_slice(s, n): return None 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) + group_size=a.group_size, bits_map=bits_map, arch=a.arch) if not out: # shard senza MTP/idx: niente file (come il download path) done[key] = "" else: @@ -749,10 +822,15 @@ def get_slice(s, n): return None # 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"]: + for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"] \ + + (["chat_template.jinja"] if a.arch == "m3" else []): src = os.path.join(a.indir, fn) - if os.path.exists(src): shutil.copy(src, a.outdir); copied.append(fn) - else: missing.append(fn) + if not os.path.exists(src): missing.append(fn); continue + if fn == "config.json" and a.arch == "m3": + write_m3_config(src, a.outdir) # flatten text_config for the engine + else: + shutil.copy(src, a.outdir) + copied.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)}" @@ -982,8 +1060,12 @@ def _download_single(url, fn, out, part, expected): if not shards: print("ERROR: no .safetensors shards found in this repository.", flush=True) return - for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"]: - try: shutil.copy(hf_hub_download(a.repo, fn, local_dir=a.outdir+"/_meta"), a.outdir) + for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"] \ + + (["chat_template.jinja"] if a.arch == "m3" else []): + try: + src = hf_hub_download(a.repo, fn, local_dir=a.outdir+"/_meta") + if fn == "config.json" and a.arch == "m3": write_m3_config(src, a.outdir) + else: shutil.copy(src, a.outdir) except Exception: pass tmp = os.path.join(a.outdir, "_inflight"); os.makedirs(tmp, exist_ok=True) if a.mtp: @@ -1035,7 +1117,7 @@ def _download_single(url, fn, out, part, expected): shutil.rmtree(tmp, ignore_errors=True); print("[IDX] DONE."); return 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, - "proj_bits": dict(PROJ_BITS)} + "proj_bits": dict(PROJ_BITS), "arch": a.arch} if not check_or_record_params(a.outdir, "out-", params): return for i, sh in enumerate(shards): if free_gb(a.outdir) < a.min_free_gb: @@ -1044,7 +1126,7 @@ def _download_single(url, fn, out, part, expected): if os.path.exists(outp): continue # gia' fatto -> ripartibile print(f"[{i+1}/{len(shards)}] downloading {sh} ({free_gb(a.outdir):.0f} GB free)...", flush=True) p = download_retry(a.repo, sh, tmp) - out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map) + out = {}; convert_shard(p, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map, arch=a.arch) _save_file_atomic(save_file, out, outp) os.remove(p) # <-- cancella subito lo shard fp8 for blob in glob.glob(os.path.join(tmp, "**", "*"), recursive=True): From 5b6db17ca1aec93df5ed4b813d0fa42496938975 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Wed, 22 Jul 2026 12:26:40 +0200 Subject: [PATCH 02/22] =?UTF-8?q?engine:=20MiniMax-M3=20architecture=20sup?= =?UTF-8?q?port=20(ARCH=5FM3)=20=E2=80=94=20GQA=20attention,=20swigluoai,?= =?UTF-8?q?=20Gemma=20norms?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit model_type minimax_* in config.json switches the engine to the M3 family: - load_cfg: M3 key set (num_local_experts, head_dim, num_key_value_heads, rotary_dim, dense/shared intermediate, swiglu params, moe_layer_freq -> first_dense); the GQA KV rows ride the MLA cache aliases (Lc=K, Rc=V, kv_lora/qk_rope = n_kv_heads*head_dim) so kv_alloc/bind/persist/mux work unchanged; arch-split config validation. - loader: self_attn.{q,k,v,o}_proj + per-head q_norm/k_norm for ARCH_M3; MoE/router/shared names already line up via the converter mapping. - attention_gqa(): per-head Gemma QK-norm BEFORE partial split-half (NEOX) RoPE on the first rotary_dim dims, standard per-token KV rows, grouped scores/softmax/values (repeat_kv h -> h/(H/NK)), o_proj; honors the ragged kvs/positions contract; full causal attention (MSA block selection exact <=2048 ctx, deferred beyond). - act_glu(): all 11 silu(g)*u sites now dispatch silu vs swigluoai (clamps, alpha, and the (up+1) factor); rmsnorm gains the Gemma (1+w) variant behind g_gemma_norm. Both default off -> GLM math untouched. K3's parallel combine moves from expert_ffn's body into act_glu, where it covers both activations and every caller; elementwise, so bit-identical. - registry: a FamilyDescriptor entry, not per-file branches. minimax_m3 (and the raw minimax_m3_vl wrapper) resolve to engine_artifact "colibri" with internal_arch "minimax_m3" -- the shared binary self-dispatches on config.json, and the descriptor keeps the two families distinct where it matters (template, planner, limits, model id). _minimax_geometry() mirrors kv_alloc's GQA rows + the MTP row and attention_gqa()'s scratch set; the registry's own minimax double is superseded by the shipped descriptor, so its test now pins the production geometry. Registering a family also obligates a segment-conformance row (release_policy "all_registered_families"), so the manifest, the C fixture matrix and its roster gain minimax_m3: GQA KV + the MSA indexer, which reuses the DSA indexer's own cache slots in the engine. The REAL segment adapter is a separate subsystem and stays at six; the manifest records that M3 needs one. - VK: expert tier + fused shared expert gated off for M3 (shaders hardcode silu); the per-matmul dense chain stays available. - chat: MiniMax template (]~!b[ / ]~b]role / [e~[) in the serve dialog loop, coli run, and openai_server (text subset; tools 400 for now); [e~[ EOS fallback. - tools: make_m3tiny.py (tiny random M3 checkpoint) + oracle_m3.py (numpy reference on the DEQUANTIZED container -> the engine's REF/TF gate). --- c/coli | 18 +- c/colibri.c | 253 +++++++++++++++++++--- c/family_registry.py | 69 ++++++ c/openai_server.py | 33 ++- c/tests/segment_conformance_fixtures.c | 15 ++ c/tests/segment_conformance_manifest.json | 7 + c/tests/test_family_registry.py | 56 ++--- c/tests/test_segment_conformance.c | 4 +- c/tools/convert_fp8_to_int4.py | 2 + c/tools/make_m3tiny.py | 70 ++++++ c/tools/oracle_m3.py | 120 ++++++++++ 11 files changed, 574 insertions(+), 73 deletions(-) create mode 100644 c/tools/make_m3tiny.py create mode 100644 c/tools/oracle_m3.py diff --git a/c/coli b/c/coli index 804d60988..e369ee4ea 100755 --- a/c/coli +++ b/c/coli @@ -1220,11 +1220,19 @@ def cmd_run(a): [engine, str(cap_for_launch(a.cap,e,16)), "8"], input=prompt+"\n", text=True, env=e, check=False) sys.exit(result.returncode) - # template ufficiale GLM-5.2: niente \n dopo i ruoli; = risposta diretta (nothink). - # THINK=1 lascia aperto, stessa convenzione del serve mode (glm.c). EN: THINK=1 leaves - # open so the engine emits its reasoning block; the default stays nothink. - tk="" if os.environ.get("THINK","0")=="1" else "" - e=env_for(a); e["PROMPT"]=f"[gMASK]<|user|>{prompt}<|assistant|>{tk}" + # Official GLM-5.2 template: no \n after the role markers; means a + # direct answer (nothink). THINK=1 leaves open so the engine emits its + # reasoning block -- same convention as serve mode (glm.c); the default stays nothink. + think=os.environ.get("THINK","0")=="1" + if arch=="minimax_m3": + # MiniMax-M3 chat_template.jinja: ]~!b[ once, ]~b]\n[e~[\n blocks, + # ai turn opens bare (adaptive thinking); prefix = nothink, like history. + tk="" if think else "" + p=f"]~!b[]~b]user\n{prompt}[e~[\n]~b]ai\n{tk}" + else: + tk="" if think else "" + p=f"[gMASK]<|user|>{prompt}<|assistant|>{tk}" + e=env_for(a); e["PROMPT"]=p sys.exit(subprocess.call([GLM, str(cap_for_launch(a.cap,e,0))], env=e)) def server_probe(base, api_key=None, timeout=1.5): diff --git a/c/colibri.c b/c/colibri.c index 4a1d8f903..4744922f5 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -160,6 +160,16 @@ static int g_cluster_n; * there -- see coli_resolve_cap() and cap_for_ram()'s CAP_RAISE default. */ static int g_ssd_fast; +/* Architecture family. ARCH_GLM = the original DeepSeek/GLM MLA + noaux_tc line the + * engine was built around. ARCH_M3 = MiniMax-M3 (minimax_m3_vl text model): standard + * GQA attention (per-head Gemma QK-norm before a split-half partial RoPE), Gemma-style + * RMSNorms (x*(1+w)), swigluoai activation, and the SAME router math as GLM (sigmoid + + * e_score_correction_bias choice, renormalized raw-sigmoid weights, routed scaling). + * For M3 the MLA cache aliases carry GQA rows: Lc = K rows, Rc = V rows, both + * n_kv_heads*head_dim wide (kv_lora/qk_rope are set to that stride at load). */ +#define ARCH_GLM 0 +#define ARCH_M3 1 + typedef struct { int hidden, n_layers, n_heads, n_experts, topk, moe_inter, dense_inter; int first_dense, q_lora, kv_lora, qk_nope, qk_rope, qk_head, v_head, n_shared, vocab; @@ -168,6 +178,10 @@ typedef struct { int index_topk, index_nh, index_hd; /* DSA lightning indexer */ int8_t idx_type[128]; /* per layer: 1=full (calcola), 0=shared (riusa) */ float eps, theta, attn_scale, routed_scale; + int arch; /* ARCH_GLM | ARCH_M3 */ + int n_kv_heads, head_dim, rotary; /* M3: GQA KV heads, head dim, partial-rope dims */ + int shared_inter; /* M3: shared-expert intermediate (GLM: moe_inter*n_shared) */ + float swiglu_alpha, swiglu_limit; /* M3: swigluoai parameters */ } Cfg; /* tensore [O,I] in uno di tre formati: @@ -365,6 +379,8 @@ typedef struct { float *in_ln, *post_ln; /* MLA (densa, quantizzata) */ QT q_a, q_b, kv_a, kv_b, o; float *q_a_ln, *kv_a_ln; + /* GQA (ARCH_M3): plain projections + per-head QK-norm weights [head_dim] */ + QT q_p, k_p, v_p; float *q_hn, *k_hn; #ifdef COLI_CUDA ColiCudaTensor *kv_b_shard[COLI_CUDA_MAX_DEVICES]; int shard_h0[COLI_CUDA_MAX_DEVICES],shard_hn[COLI_CUDA_MAX_DEVICES],n_kv_b_shard; @@ -1472,9 +1488,16 @@ static void qt_fill(QT *t, const float *w, int bits){ else pack_int4(w, t->q4, t->s, t->O, t->I, bits); } +/* ARCH_M3 numeric conventions, set once at model_init from the config. Both default + * to the GLM behavior so every existing path is bit-identical when they stay 0. */ +static int g_gemma_norm=0; /* rmsnorm scales by (1+w) instead of w */ +static int g_act_swigluoai=0; /* glu = clamp+alpha-sigmoid, (up+1)*glu */ +static float g_swiglu_alpha=1.702f, g_swiglu_limit=7.0f; static void rmsnorm(float *out, const float *x, const float *w, int D, float eps){ double ms=0; for(int i=0;i=16384 && !omp_in_parallel()) + for(int64_t i=0;i=16384 && !omp_in_parallel()) + for(int64_t i=0;iL?L:u[i]); + g[i]=(uv+1.f)*(gv/(1.f+expf(-A*gv))); + } +} /* FFN di un expert MoE (routed): gate+up -> silu(gate)*up -> down. Un solo helper per * i ~6 siti inline di moe() e (piu' avanti) il worker distribuito, che la chiama con i @@ -1506,13 +1553,7 @@ static inline float siluf(float x){ return x/(1.f+expf(-x)); } * expert, vedi il commento MB_BUILD) e' la correzione di un bug latente, NON un no-op. */ static void expert_ffn(float *hh, float *gg, float *uu, const float *xg, QT *g, QT *u, QT *d, int nr, int I){ expert_gate_up(gg,uu,xg,g,u,nr); - /* K3: silu*up in parallelo — per-elemento, nessun ordine condiviso: gli - * stessi bit in qualunque schedulazione. Sotto soglia resta seriale (il - * fork/join costerebbe piu' del loop). - * EN: elementwise silu*up, parallel above a size threshold; per-element - * math has no shared accumulation order, so bits are unchanged. */ - #pragma omp parallel for schedule(static) if((int64_t)nr*I>=16384) - for(int64_t z=0;z<(int64_t)nr*I;z++) gg[z]=siluf(gg[z])*uu[z]; + act_glu(gg, uu, (int64_t)nr*I); /* silu(g)*u, or swigluoai when the arch sets it */ if(d->fmt==6) e8_rot_rows(gg,nr,I); /* down input is per-expert — rotate here */ #if !defined(COLI_CUDA)&&!defined(COLI_METAL)&&!defined(COLI_VULKAN) /* K3: la meta' mancante di #1071 — la 0.1 sollevo' la quantizzazione della @@ -1625,6 +1666,8 @@ static jval* cfg_root(const char *snap, char **arena){ static int gi(jval*r,const char*k){ jval*v=json_get(r,k); return v?(int)v->num:0; } static void load_cfg(Cfg *c, const char *snap){ char *ar=NULL; jval *r=cfg_root(snap,&ar); + { jval *mt=json_get(r,"model_type"); + c->arch = (mt && mt->str && !strncmp(mt->str,"minimax",7)) ? ARCH_M3 : ARCH_GLM; } c->hidden=gi(r,"hidden_size"); c->n_layers=gi(r,"num_hidden_layers"); c->n_heads=gi(r,"num_attention_heads"); c->n_experts=gi(r,"n_routed_experts"); c->topk=gi(r,"num_experts_per_tok"); c->moe_inter=gi(r,"moe_intermediate_size"); @@ -1637,7 +1680,32 @@ static void load_cfg(Cfg *c, const char *snap){ jval *ep=json_get(r,"rms_norm_eps"); c->eps=ep?(float)ep->num:1e-5f; jval *rs=json_get(r,"routed_scaling_factor"); c->routed_scale=rs?(float)rs->num:1.f; jval *rp=json_get(r,"rope_parameters"); jval *th=rp?json_get(rp,"rope_theta"):NULL; + if(!th) th=json_get(r,"rope_theta"); /* M3 (and some GLM configs): flat key */ c->theta = th?(float)th->num:10000.f; + if(c->arch==ARCH_M3){ + /* MiniMax-M3 key names (flattened text_config, see tools/convert --arch m3). + * intermediate_size IS the routed-expert width there; the first-3-dense + * pattern arrives as a moe_layer_freq 0/1 array (count the leading zeros). */ + c->n_experts=gi(r,"num_local_experts"); + c->moe_inter=gi(r,"intermediate_size"); + c->dense_inter=gi(r,"dense_intermediate_size"); + c->shared_inter=gi(r,"shared_intermediate_size"); + if(!c->shared_inter) c->shared_inter=c->moe_inter*(c->n_shared?c->n_shared:1); + c->n_kv_heads=gi(r,"num_key_value_heads"); + c->head_dim=gi(r,"head_dim"); + c->rotary=gi(r,"rotary_dim"); + { jval *sa=json_get(r,"swiglu_alpha"); c->swiglu_alpha=sa?(float)sa->num:1.702f; + jval *sl=json_get(r,"swiglu_limit"); c->swiglu_limit=sl?(float)sl->num:7.0f; } + { jval *fr=json_get(r,"moe_layer_freq"); c->first_dense=0; + if(fr && fr->t==J_ARR){ int i=0; while(ilen && (int)fr->kids[i]->num==0) i++; + c->first_dense=i; } } + c->norm_topk=1; /* M3 always renormalizes the top-k weights */ + c->n_group=1; /* no grouped routing (key absent -> 0) */ + /* GQA rows ride the MLA cache aliases: Lc rows = K (n_kv*hd), Rc rows = V. */ + c->kv_lora = c->n_kv_heads*c->head_dim; + c->qk_rope = c->n_kv_heads*c->head_dim; + c->qk_nope=0; c->v_head=c->head_dim; c->q_lora=0; + } /* token di stop: GLM-5.2 ne ha TRE (endoftext, user, observation). Fermarsi solo sul * primo = generare spazzatura invisibile dopo la fine del turno (5-10x token sprecati). */ c->n_stop=0; @@ -1685,8 +1753,13 @@ static void load_cfg(Cfg *c, const char *snap){ c->idx_type[i] = !strcmp(it->kids[i]->str,"full"); else { int v=i-off+1; if(v<0) v=0; c->idx_type[i] = (v%freq)==0; } } } - c->qk_head=c->qk_nope+c->qk_rope; - c->attn_scale = 1.f / sqrtf((float)c->qk_head); + if(c->arch==ARCH_M3){ + c->qk_head=c->head_dim; + c->attn_scale = 1.f / sqrtf((float)c->head_dim); + } else { + c->qk_head=c->qk_nope+c->qk_rope; + c->attn_scale = 1.f / sqrtf((float)c->qk_head); + } if(c->n_group!=1){ fprintf(stderr,"this engine requires n_group=1 (GLM-5.2)\n"); exit(1); } /* VALIDAZIONE (report PR #25): il config.json arriva da mirror non fidati — dimensioni * ostili non devono superare questo punto. Un solo choke point protegge ogni alloc a valle. */ @@ -1696,11 +1769,18 @@ static void load_cfg(Cfg *c, const char *snap){ CKR("num_attention_heads",c->n_heads,1,1024) CKR("n_routed_experts",c->n_experts,1,4096) CKR("num_experts_per_tok",c->topk,1,64) CKR("moe_intermediate_size",c->moe_inter,1,1<<20) CKR("intermediate_size",c->dense_inter,1,1<<24) CKR("first_k_dense_replace",c->first_dense,0,c->n_layers) + CKR("n_shared_experts",c->n_shared,0,64) CKR("vocab_size",c->vocab,1,1<<24) + if(c->arch==ARCH_M3){ + CKR("num_key_value_heads",c->n_kv_heads,1,c->n_heads) + CKR("head_dim",c->head_dim,1,1<<16) CKR("rotary_dim",c->rotary,0,c->head_dim) + if(c->n_heads % c->n_kv_heads){ fprintf(stderr,"config: num_attention_heads %% num_key_value_heads != 0\n"); exit(1); } + if(c->rotary % 2){ fprintf(stderr,"config: rotary_dim must be even\n"); exit(1); } + } else { CKR("q_lora_rank",c->q_lora,0,1<<20) CKR("kv_lora_rank",c->kv_lora,1,1<<20) CKR("qk_nope_head_dim",c->qk_nope,1,1<<16) CKR("qk_rope_head_dim",c->qk_rope,1,1<<16) - CKR("v_head_dim",c->v_head,1,1<<16) CKR("n_shared_experts",c->n_shared,0,64) - CKR("vocab_size",c->vocab,1,1<<24) CKR("index_topk",c->index_topk,0,1<<20) + CKR("v_head_dim",c->v_head,1,1<<16) CKR("index_topk",c->index_topk,0,1<<20) CKR("index_n_heads",c->index_nh,0,1024) CKR("index_head_dim",c->index_hd,0,1<<16) + } if(c->topk>c->n_experts){ fprintf(stderr,"config: num_experts_per_tok=%d exceeds n_routed_experts=%d\n", c->topk,c->n_experts); exit(1); } @@ -2256,6 +2336,13 @@ static void model_init_range(Model *m, const char *snap, int cap, int init_telemetry){ memset(m,0,sizeof(*m)); m->ebits=ebits; m->dbits=dbits; load_cfg(&m->c,snap); + if(m->c.arch==ARCH_M3){ + g_gemma_norm=1; g_act_swigluoai=1; + g_swiglu_alpha=m->c.swiglu_alpha; g_swiglu_limit=m->c.swiglu_limit; + fprintf(stderr,"[ARCH] MiniMax-M3: GQA %d/%d heads hd %d, rotary %d, %d experts top-%d, " + "gemma-norm + swigluoai\n", m->c.n_heads, m->c.n_kv_heads, m->c.head_dim, + m->c.rotary, m->c.n_experts, m->c.topk); + } { const char *xd=getenv("COLI_MODEL_DIRS"); /* SPLIT: model shards spread across N drives */ st_init_multi(&m->S,snap,(xd&&*xd)?xd:NULL); } Cfg *c=&m->c; char nm[256]; int H=c->n_heads, D=c->hidden; @@ -2311,6 +2398,14 @@ static void model_init_range(Model *m, const char *snap, int cap, #define P(s) (snprintf(nm,sizeof(nm),"model.layers.%d." s,i),nm) l->in_ln=ld(m,P("input_layernorm.weight")); l->post_ln=ld(m,P("post_attention_layernorm.weight")); + if(c->arch==ARCH_M3){ + l->q_p = qt_load(m,P("self_attn.q_proj.weight"), H*c->head_dim, D, dbits); + l->k_p = qt_load(m,P("self_attn.k_proj.weight"), c->n_kv_heads*c->head_dim, D, dbits); + l->v_p = qt_load(m,P("self_attn.v_proj.weight"), c->n_kv_heads*c->head_dim, D, dbits); + l->o = qt_load(m,P("self_attn.o_proj.weight"), D, H*c->head_dim, dbits); + l->q_hn = ld(m,P("self_attn.q_norm.weight")); + l->k_hn = ld(m,P("self_attn.k_norm.weight")); + } else { l->q_a = qt_load(m,P("self_attn.q_a_proj.weight"), c->q_lora, D, dbits); l->q_a_ln= ld(m,P("self_attn.q_a_layernorm.weight")); l->q_b = qt_load(m,P("self_attn.q_b_proj.weight"), H*c->qk_head, c->q_lora, dbits); @@ -2318,6 +2413,7 @@ static void model_init_range(Model *m, const char *snap, int cap, l->kv_a_ln= ld(m,P("self_attn.kv_a_layernorm.weight")); l->kv_b = qt_load(m,P("self_attn.kv_b_proj.weight"), H*(c->qk_nope+c->v_head), c->kv_lora, dbits); l->o = qt_load(m,P("self_attn.o_proj.weight"), D, H*c->v_head, dbits); + } #ifdef COLI_CUDA qt_cuda_colocate(&l->o,&l->kv_b); qt_cuda_colocate(&l->q_a,&l->kv_b); /* PIPE: intera catena attention sulla */ @@ -2335,7 +2431,7 @@ static void model_init_range(Model *m, const char *snap, int cap, } else { l->router=ld(m,P("mlp.gate.weight")); l->router_bias=ld(m,P("mlp.gate.e_score_correction_bias")); - int sI=c->moe_inter*c->n_shared; + int sI = c->arch==ARCH_M3 ? c->shared_inter : c->moe_inter*c->n_shared; l->sh_gate = qt_load(m,P("mlp.shared_experts.gate_proj.weight"), sI, D, dbits); l->sh_up = qt_load(m,P("mlp.shared_experts.up_proj.weight"), sI, D, dbits); l->sh_down = qt_load(m,P("mlp.shared_experts.down_proj.weight"), D, sI, dbits); @@ -4084,10 +4180,99 @@ static void kv_lc_rows_f32(Model *m, int layer, int64_t t0, int64_t n, float *ds dst+(t-t0)*c->kv_lora, c->kv_lora); } } +/* ---------------- ARCH_M3: standard GQA attention -------------------------------- + * MiniMax-M3 text backbone (transformers modeling_minimax_m3_vl conventions): + * per-head Gemma RMSNorm (x*(1+w), eps=rms_norm_eps) on Q and K BEFORE RoPE; partial + * split-half (NEOX) RoPE on the first c->rotary dims of each head; standard per-token + * KV rows in the Lc/Rc aliases (K = Lc, V = Rc, both n_kv_heads*head_dim wide — set + * up by load_cfg); score/softmax/value with H/n_kv_heads query heads per KV head; + * o_proj at the end. Mirrors attention_rows' ragged contract (kvs/positions per-row + * KV target, else pos_base+s). MSA block selection is NOT implemented: full causal + * attention — EXACT for windows up to sparse_topk_blocks*block = 2048 tokens (the + * indexer selects every block then), an approximation beyond. */ +static void rope_half_neox(float *v, int rot, int pos, float theta){ + int h2=rot/2; + for(int j=0;jc; int H=c->n_heads, NK=c->n_kv_heads, hd=c->head_dim; + int G=H/NK, KVd=NK*hd, rot=c->rotary; + double ta0=now_s(); + float *q=falloc((int64_t)S*H*hd), *k=falloc((int64_t)S*KVd), *v=falloc((int64_t)S*KVd); + float *ctx=falloc((int64_t)S*H*hd); + double tp0=now_s(); + matmul_qt(q,x,&l->q_p,S); + matmul_qt(k,x,&l->k_p,S); + matmul_qt(v,x,&l->v_p,S); + m->t_aproj+=now_s()-tp0; + for(int s=0;skv; + int pos=positions?positions[s]:pos_base+s; + float *qs=q+(int64_t)s*H*hd, *ksr=k+(int64_t)s*KVd, *vsr=v+(int64_t)s*KVd; + for(int h=0;hq_hn,hd,c->eps); rope_half_neox(qh,rot,pos,c->theta); } + for(int h=0;hk_hn,hd,c->eps); rope_half_neox(kh,rot,pos,c->theta); } + memcpy(ks->Lc[layer]+(int64_t)pos*KVd, ksr, (size_t)KVd*sizeof(float)); + memcpy(ks->Rc[layer]+(int64_t)pos*KVd, vsr, (size_t)KVd*sizeof(float)); + } + double tc0=now_s(); + for(int s=0;skv; + int pos=positions?positions[s]:pos_base+s; + int st0=ks->kv_start[layer], T=pos+1, win=T-st0; + const float *K0=ks->Lc[layer], *V0=ks->Rc[layer]; + float *qs=q+(int64_t)s*H*hd, *cs=ctx+(int64_t)s*H*hd; + float *att=falloc((int64_t)H*win); + #pragma omp parallel for schedule(static) + for(int h=0;hattn_scale; ah[t-st0]=d; if(d>mx) mx=d; + } + float ssum=0; + for(int t=0;tt_acore+=now_s()-tc0; + double to0=now_s(); + matmul_qt(out,ctx,&l->o,S); + m->t_aout+=now_s()-to0; + free(q); free(k); free(v); free(ctx); + m->t_attn += now_s()-ta0; +} + static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int pos_base, KVState *const *kvs, const int *positions, float *out){ Cfg *c=&m->c; int H=c->n_heads, D=c->hidden, qh=c->qk_head, vh=c->v_head; int kvb_dim=H*(c->qk_nope+vh), Tk=pos_base+S; + if(c->arch==ARCH_M3){ attention_gqa(m,l,layer,x,S,pos_base,kvs,positions,out); return; } double ta0=now_s(); #ifdef COLI_METAL /* Fused decode attention on GPU: whole layer in one command buffer (keeps the GPU hot). @@ -5778,7 +5963,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int float *gj=GG+(int64_t)j*I, *uj=UU+(int64_t)j*I; for(int o=c0;og.s[o]*sx0; for(int o=c0;ou.s[o]*sx0; - for(int o=c0;osh_gate.fmt; - if(g_vk_dense && !omp_in_parallel() && (fsh==1||fsh==2||fsh==5) && - l->sh_up.fmt==fsh && l->sh_down.fmt==fsh){ + if(g_vk_dense && !g_act_swigluoai && !omp_in_parallel() && (fsh==1||fsh==2||fsh==5) && + l->sh_up.fmt==fsh && l->sh_down.fmt==fsh){ /* fused shader hardcodes silu — M3 uses the per-matmul chain */ #define SW_(t) ((t).fmt==1?(const void*)(t).q8:(const void*)(t).q4) if(coli_vk_tensor_ensure(&l->sh_gate.vk,SW_(l->sh_gate),l->sh_gate.s,fsh,D,sI,l->sh_gate.gs)&& coli_vk_tensor_ensure(&l->sh_up.vk, SW_(l->sh_up), l->sh_up.s, fsh,D,sI,l->sh_up.gs)&& @@ -6214,7 +6399,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int if(!vk_matmul_qt(&l->sh_up, su, x, S)) #endif matmul_qt(su, x, &l->sh_up, S); - for(int64_t z=0;z<(int64_t)S*sI;z++) sg[z]=siluf(sg[z])*su[z]; + act_glu(sg, su, (int64_t)S*sI); #ifdef COLI_VULKAN if(!vk_matmul_qt(&l->sh_down, hh, sg, S)) #endif @@ -6246,7 +6431,7 @@ static void dense_mlp(Layer *l, float *x, int S, int D, int I, float *out){ float *g=falloc((int64_t)S*I), *u=falloc((int64_t)S*I); matmul_qt(g, x, &l->gate_proj, S); matmul_qt(u, x, &l->up_proj, S); - for(int64_t i=0;i<(int64_t)S*I;i++) g[i]=siluf(g[i])*u[i]; + act_glu(g, u, (int64_t)S*I); matmul_qt(out, g, &l->down_proj, S); free(g); free(u); } @@ -6282,7 +6467,7 @@ static void la_predict(Model *m, int target, const float *h, int kind){ rmsnorm(snrm, h, sl->post_ln, D, c->eps); matmul_qt(sg, snrm, &sl->sh_gate, 1); matmul_qt(su, snrm, &sl->sh_up, 1); - for(int i=0;ish_down, 1); for(int i=0;ipost_ln, D, c->eps); @@ -6630,7 +6815,7 @@ static void pilot_prefetch(Model *m, int lnext, const float *x, int S){ rmsnorm(snrm, xs, sl->post_ln, D, c->eps); matmul_qt(sg, snrm, &sl->sh_gate, 1); matmul_qt(su, snrm, &sl->sh_up, 1); - for(int i=0;ish_down, 1); for(int i=0;ipost_ln, D, c->eps); @@ -9015,6 +9200,7 @@ static void run_serve(Model *m, const char *snap){ char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); Tok T; tok_load(&T,tkp); int eos=tok_id_of(&T,"<|endoftext|>"); + if(eos<0) eos=tok_id_of(&T,"[e~["); /* MiniMax-M3 end-of-sequence */ stops_arm_tok(&m->c, eos, &T); grammar_setup(&g_grd,&T); /* metodo F: GRAMMAR=file.gbnf (#48) */ if(g_temp<0) g_temp=0.7f; /* auto: 0.7, NON l'1.0 ufficiale — la coda della @@ -9091,8 +9277,13 @@ static void run_serve(Model *m, const char *snap){ int bl=0, k=0; /* costruisce/tokenizza il turno */ /* template UFFICIALE GLM-5.2 (chat_template.jinja): niente \n dopo i ruoli, e dopo * <|assistant|> serve SEMPRE il blocco think — lo DISATTIVA (nothink): - * col template sbagliato il modello farfuglia e non emette mai lo stop. THINK=1 lo abilita. */ - const char *tk = getenv("THINK")&&atoi(getenv("THINK"))? "" : ""; + * col template sbagliato il modello farfuglia e non emette mai lo stop. THINK=1 lo abilita. + * MiniMax-M3 (chat_template.jinja): ]~!b[ once, then ]~b]\n[e~[\n blocks + * with roles user/ai; the ai turn opens bare (adaptive thinking) — THINK=0-style + * suppression prefixes the reply with , mirroring the official history form. */ + int think_on = getenv("THINK")&&atoi(getenv("THINK")); + const char *tk = m->c.arch==ARCH_M3 ? (think_on? "" : "") + : (think_on? "" : ""); if(raw_mode){ int *tmp=malloc(maxctx*sizeof(int)); if(!tmp){fprintf(stderr,"OOM raw tokens\n");exit(1);} prompt_tokens=tok_encode(&T,input,input_n,tmp,maxctx-8-g_draft); @@ -9109,12 +9300,19 @@ static void run_serve(Model *m, const char *snap){ active,len,prompt_tokens,k); free(tmp); } else { - if(templ){ if(first) bl+=snprintf(buf+bl,(1<<16)-bl,"[gMASK]"); - bl+=snprintf(buf+bl,(1<<16)-bl,"<|user|>%s<|assistant|>%s",input,tk); } + if(templ){ if(m->c.arch==ARCH_M3){ + if(first) bl+=snprintf(buf+bl,(1<<16)-bl,"]~!b["); + bl+=snprintf(buf+bl,(1<<16)-bl,"]~b]user\n%s[e~[\n]~b]ai\n%s",input,tk); + } else { + if(first) bl+=snprintf(buf+bl,(1<<16)-bl,"[gMASK]"); + bl+=snprintf(buf+bl,(1<<16)-bl,"<|user|>%s<|assistant|>%s",input,tk); + } } else bl+=snprintf(buf+bl,(1<<16)-bl,"%s",input); k=tok_encode(&T,buf,bl,hist+len,maxctx-len); prompt_tokens=k; if(len+k+8+g_draft>=maxctx){ len=0; first=1; kv_disk_reset(m); - bl=0; if(templ){ bl+=snprintf(buf+bl,(1<<16)-bl,"[gMASK]<|user|>%s<|assistant|>%s",input,tk); } + bl=0; if(templ){ bl+=snprintf(buf+bl,(1<<16)-bl, m->c.arch==ARCH_M3 + ? "]~!b[]~b]user\n%s[e~[\n]~b]ai\n%s" + : "[gMASK]<|user|>%s<|assistant|>%s",input,tk); } else bl+=snprintf(buf+bl,(1<<16)-bl,"%s",input); k=tok_encode(&T,buf,bl,hist,maxctx); if(k>maxctx-8-g_draft) k=maxctx-8-g_draft; prompt_tokens=k; @@ -9230,6 +9428,11 @@ static void vk_dense_preload(Model *m){ static void vk_registry_fill(Model *m){ Cfg *c=&m->c; int E=c->n_experts, NL=c->n_layers; if(!g_vulkan || g_vk_budget<=0) return; + if(c->arch!=ARCH_GLM){ /* expert-group shaders hardcode silu */ + fprintf(stderr,"[VK] expert tier disabled for this architecture (shader activation " + "is silu; swigluoai shader variant pending) — experts run on the CPU\n"); + return; + } int64_t nz=0; for(int i=0;ieusage[i]) for(int e=0;eeusage[i][e]) nz++; if(!nz){ fprintf(stderr,"[VK] expert tier: no usage history yet — tier empty this run " diff --git a/c/family_registry.py b/c/family_registry.py index ca9d3a1f0..951a5323e 100644 --- a/c/family_registry.py +++ b/c/family_registry.py @@ -393,6 +393,38 @@ def _dsv4_geometry(config, context, _model_dir): return PlannerGeometry(state, fixed, workspace, experts) +def _minimax_geometry(config, context, _model_dir): + """MiniMax-M3: plain GQA -- no latent compression, so the cache holds the + full K and V rows rather than a compressed pair. + + Mirrors colibri.c's ARCH_M3 allocation. load_cfg aliases the MLA cache + fields onto the GQA widths (kv_lora = qk_rope = num_key_value_heads * + head_dim, colibri.c:1411-1412), so kv_alloc's per-layer Lc/Rc pair is one + full K row and one full V row per token. It keeps num_hidden_layers + 1 + rows: the extra one is the MTP layer's KV. + + state = (layers + 1) * context * 2 * kv_heads * head_dim * 4 + + Workspace mirrors attention_gqa()'s scratch set with S == context at + prefill: q and ctx are heads-wide, k and v are kv_heads-wide, plus the + per-row score buffer (heads * window, window <= context). + + ws = context * ((2 * heads + 2 * kv_heads) * head_dim + heads) * 4 + + Experts: configured_experts = num_local_experts (M3's spelling of + n_routed_experts). + """ + layers = _required_int(config, "num_hidden_layers", "minimax_m3") + experts = _required_int(config, "num_local_experts", "minimax_m3") + heads = _required_int(config, "num_attention_heads", "minimax_m3") + kv_heads = _required_int(config, "num_key_value_heads", "minimax_m3") + head_dim = _required_int(config, "head_dim", "minimax_m3") + + state = (layers + 1) * context * 2 * kv_heads * head_dim * 4 + workspace = context * ((2 * heads + 2 * kv_heads) * head_dim + heads) * 4 + return PlannerGeometry(state, 0, workspace, experts) + + _GLM_EXPERT = re.compile( r"(?:^|\.)model\.layers\.(\d+)\.mlp\.experts\.(\d+)\." ) @@ -589,6 +621,43 @@ def _inkling_expert_inventory(name, size, config): has_gateway_adapter=True, has_cli_adapter=True, ), + FamilyDescriptor( + id="minimax_m3", + # The converter flattens text_config to the root and stamps + # "minimax_m3"; a raw MiniMax checkpoint is the VL wrapper + # ("minimax_m3_vl") with the text model nested. Both resolve here, and + # config_section="text_config" reads either shape. + model_types=("minimax_m3", "minimax_m3_vl"), + display_name="MiniMax-M3", + display_scale="426B", + # Shares GLM's binary ON PURPOSE: colibri.c reads config.json and + # switches itself to ARCH_M3 (GQA + MSA). The descriptor keeps them + # distinct where it matters -- internal_arch, template, planner -- so + # this is routing, not the #879 fallthrough. engine_aliases mirrors + # GLM's because it is literally the same artifact on disk. + engine_artifact="colibri", + engine_aliases=("glm",), + engine_group="colibri-core", + internal_arch="minimax_m3", + build_target="colibri", + process_names=("colibri",), + default_model_id="minimax-m3-colibri", + cli_adapter="minimax_m3", + gateway_adapter="minimax_m3", + planner_id="minimax_m3_gqa", + planner_geometry=_minimax_geometry, + planner_unsupported_reason="", + expert_inventory=_individual_expert_inventory(_GLM_EXPERT), + config_section="text_config", + # implicit_cap 0 is GLM's platform-auto sentinel, which this binary + # resolves itself -- the legacy 8 belongs to the sister engines. + # One KV slot: batched serve is not validated for M3 in this PR. + limits=FamilyLimits(4096, 1048576, 1024, 16384, 1, 0, "CTX"), + capabilities=FamilyCapabilities(False, False, False, True), + has_gateway_adapter=True, + has_cli_adapter=True, + tune_prompt_template="]~!b[]~b]user\n{prompt}[e~[\n]~b]ai\n", + ), ) diff --git a/c/openai_server.py b/c/openai_server.py index 1617f994f..1ec8acbd9 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -1274,6 +1274,36 @@ def render_chat_inkling(messages, enable_thinking=False, reasoning_effort=None, return "".join(prompt) +def render_chat_m3(messages, enable_thinking=False, reasoning_effort=None, tools=None, + tool_choice=None): + """Text-only subset of the MiniMax-M3 chat template (chat_template.jinja): + ]~!b[ once, then ]~b]\n[e~[\n blocks. A client "system" message + maps to the `developer` role (the official template reserves `system` for the + auto-injected model-identity block, which we do not fabricate here). History + assistant turns carry the prefix; the open ai turn does too unless + thinking is enabled. Tool calls: not yet rendered for this family.""" + if tools: + raise APIError(400, "tools are not yet supported for the MiniMax-M3 template.", "tools") + prompt = ["]~!b["] + for index, message in enumerate(messages): + if not isinstance(message, dict): + raise APIError(400, "Each message must be an object.", f"messages.{index}") + role = message.get("role") + raw = message.get("content") + text = content_text(raw, f"messages.{index}.content") if raw is not None else "" + if role in ("system", "developer"): + prompt.append(f"]~b]developer\n{text}[e~[\n") + elif role == "user": + prompt.append(f"]~b]user\n{text}[e~[\n") + elif role == "assistant": + prompt.append(f"]~b]ai\n{text.strip()}[e~[\n") + else: + raise APIError(400, f"Unsupported message role for MiniMax-M3: {role!r}.", + f"messages.{index}.role", "unsupported_role") + prompt.append("]~b]ai\n" + ("" if enable_thinking else "")) + return "".join(prompt) + + def render_chat(messages, enable_thinking=False, reasoning_effort=None, tools=None, tool_choice=None): """Render the text-only subset of the official GLM-5.2 chat template.""" @@ -1378,7 +1408,8 @@ def render_chat_for_arch(messages, enable_thinking=False, reasoning_effort=None, renderer = (render_chat_kimi if ARCH == "kimi" else render_chat_qwen if ARCH == "qwen36" else render_chat_v4 if ARCH == "deepseek_v4" else - render_chat_olmoe if ARCH == "olmoe" else render_chat) + render_chat_olmoe if ARCH == "olmoe" else + render_chat_m3 if ARCH == "minimax_m3" else render_chat) return renderer(messages, enable_thinking, reasoning_effort, tools, tool_choice) diff --git a/c/tests/segment_conformance_fixtures.c b/c/tests/segment_conformance_fixtures.c index 0dc87a2f9..af09ed224 100644 --- a/c/tests/segment_conformance_fixtures.c +++ b/c/tests/segment_conformance_fixtures.c @@ -91,6 +91,19 @@ static const ColiSegmentConformanceFixture g_fixtures[] = { COLI_SEGMENT_FIXTURE_DEVICE_CACHE, 5, 10, 4, 64, UINT32_C(0x44535634), }, + { + /* M3 keeps full K/V rows (no latent), and its MSA Lightning Indexer + * reuses the very cache slots the DSA indexer uses in the engine + * (index_hd / idx_type / Ic), so DSA_INDEXER is the honest kind here + * rather than a new one. Device cache: the Vulkan expert tier. */ + "minimax_m3", "MiniMax-M3", "fixture/minimax-m3-gqa-msa-v1", + "GQA key/value rows + MSA block-sparse indexer + device cache", + "tools/make_m3tiny.py", + COLI_SEGMENT_FIXTURE_KV | + COLI_SEGMENT_FIXTURE_DSA_INDEXER | + COLI_SEGMENT_FIXTURE_DEVICE_CACHE, + 3, 8, 4, 64, UINT32_C(0x4d4e4d33), + }, }; static int fail(char *error, size_t error_size, const char *message) { @@ -339,6 +352,7 @@ DECLARE_OPEN_WRAPPER(kimi, 2) DECLARE_OPEN_WRAPPER(olmoe, 3) DECLARE_OPEN_WRAPPER(qwen36, 4) DECLARE_OPEN_WRAPPER(deepseek_v4, 5) +DECLARE_OPEN_WRAPPER(minimax_m3, 6) #define FIXTURE_ADAPTER(name) \ { \ @@ -362,6 +376,7 @@ static const ColiSegmentAdapter g_adapters[] = { FIXTURE_ADAPTER(olmoe), FIXTURE_ADAPTER(qwen36), FIXTURE_ADAPTER(deepseek_v4), + FIXTURE_ADAPTER(minimax_m3), }; int coli_segment_conformance_register_fixtures(void) { diff --git a/c/tests/segment_conformance_manifest.json b/c/tests/segment_conformance_manifest.json index 6cd2f88d3..1a1d69725 100644 --- a/c/tests/segment_conformance_manifest.json +++ b/c/tests/segment_conformance_manifest.json @@ -43,6 +43,13 @@ "state_components": ["mhc", "window_attention", "compressed_attention", "compressor", "indexer"], "oracle": {"kind": "generated_tiny", "path": "tools/make_deepseek_v4_tiny.py"}, "real_adapter_required": true + }, + { + "family_id": "minimax_m3", + "state_schema": "fixture/minimax-m3-gqa-msa-v1", + "state_components": ["gqa_kv", "rope", "msa_indexer", "device_cache"], + "oracle": {"kind": "generated_tiny", "path": "tools/make_m3tiny.py"}, + "real_adapter_required": true } ] } diff --git a/c/tests/test_family_registry.py b/c/tests/test_family_registry.py index 19caa9b62..7d763068b 100644 --- a/c/tests/test_family_registry.py +++ b/c/tests/test_family_registry.py @@ -38,15 +38,6 @@ def qwen_geometry(config, context, _model_dir): return PlannerGeometry(kv, fixed, 0, config["num_experts"]) -def minimax_geometry(config, context, _model_dir): - state = ((config["num_hidden_layers"] + 1) * context * - config["num_key_value_heads"] * config["head_dim"] * 2 * 4) - sparse = config["sparse_attention_config"] - state += sum(bool(value) for value in sparse["sparse_attention_freq"]) * \ - context * sparse["sparse_index_dim"] * 4 - return PlannerGeometry(state, 0, 0, config["num_local_experts"]) - - TEST_INVENTORY = lambda _name, _size, _config: () QWEN36_FIXTURE = FamilyDescriptor( id="qwen36", @@ -70,29 +61,6 @@ def minimax_geometry(config, context, _model_dir): limits=FamilyLimits(8192, 262144, 1024, 8192, 1, 8, "Q36_MAXT"), capabilities=FamilyCapabilities(False, False, False, True), ) -MINIMAX_M3_FIXTURE = FamilyDescriptor( - id="minimax_m3", - model_types=("minimax_m3",), - display_name="MiniMax M3", - display_scale="", - engine_artifact="colibri", - engine_aliases=(), - engine_group="colibri-core", - internal_arch="minimax_m3", - build_target="colibri", - process_names=("colibri",), - default_model_id="minimax-m3-colibri", - cli_adapter="minimax_m3", - gateway_adapter="minimax_m3", - planner_id="minimax_m3_gqa", - planner_geometry=minimax_geometry, - planner_unsupported_reason="", - expert_inventory=TEST_INVENTORY, - config_section="root", - limits=FamilyLimits(8192, 262144, 1024, 8192, 1, 8, "CTX"), - capabilities=FamilyCapabilities(True, False, False, True), -) - class FamilyRegistryTest(unittest.TestCase): def test_production_descriptors_are_complete_unique_and_serializable(self): @@ -151,7 +119,7 @@ def test_qwen_fixture_models_gqa_and_fixed_deltanet_state(self): for model_type in ("qwen2", "qwen3_moe", "my_qwen_model"): self.assertNotIn(model_type, by_type) - def test_minimax_fixture_can_share_colibri_without_becoming_glm(self): + def test_minimax_shares_colibri_without_becoming_glm(self): config = { "model_type": "minimax_m3", "num_hidden_layers": 2, @@ -160,14 +128,15 @@ def test_minimax_fixture_can_share_colibri_without_becoming_glm(self): "head_dim": 8, "num_local_experts": 4, "num_experts_per_tok": 2, - "sparse_attention_config": { - "use_sparse_attention": True, - "sparse_index_dim": 8, - "sparse_attention_freq": [0, 1], - }, + "moe_layer_freq": [0, 1], } - by_id, by_type = _build_registry(FAMILIES + (MINIMAX_M3_FIXTURE,)) + # minimax_m3 is a registered family now, so the fixture this test + # landed with would collide on its id. Assert against the production + # descriptor instead -- the stronger test: it pins the shipped + # geometry, not a copy. + by_id, by_type = _build_registry(FAMILIES) family = by_type[config["model_type"]] + self.assertEqual(family, by_id["minimax_m3"]) self.assertEqual(family.engine_artifact, by_id["glm"].engine_artifact) self.assertEqual(family.engine_group, by_id["glm"].engine_group) self.assertNotEqual(family.internal_arch, by_id["glm"].internal_arch) @@ -175,7 +144,9 @@ def test_minimax_fixture_can_share_colibri_without_becoming_glm(self): "model_dir": "."})() geometry = planner_geometry(resolved, 32) self.assertEqual(geometry.configured_experts, 4) - self.assertEqual(geometry.context_state_bytes, 13_312) + # 3 rows (2 layers + the MTP row) x 32 tokens x (K + V) x 2 kv heads + # x 8 head_dim x 4 bytes. No latent compression: GQA caches full rows. + self.assertEqual(geometry.context_state_bytes, 12_288) def test_olmoe_fixture_models_conventional_fp32_kv_cache(self): # OLMoE keeps a full K and V cache per layer, sized at num_attention_heads @@ -843,6 +814,11 @@ def test_tuning_replay_prompts_are_registry_owned(self): "qwen36": "<|im_start|>user\nhello {world}<|im_end|>\n" "<|im_start|>assistant\n\n", "deepseek_v4": "hello {world}", + # MiniMax-M3 opens the ai turn with , the same nothink + # marker its history turns carry -- the bare "]~b]ai\\n" state is + # the thinking one, so replaying without it would tune a different + # workload than the default `coli run` / `coli serve` path. + "minimax_m3": "]~!b[]~b]user\nhello {world}[e~[\n]~b]ai\n", } self.assertEqual( {family.id: tuning_replay_prompt(family, prompt) for family in FAMILIES}, diff --git a/c/tests/test_segment_conformance.c b/c/tests/test_segment_conformance.c index 491954613..8ef0d5a06 100644 --- a/c/tests/test_segment_conformance.c +++ b/c/tests/test_segment_conformance.c @@ -5,7 +5,7 @@ static int expected_family(const char *family_id) { static const char *const expected[] = { - "glm", "inkling", "kimi", "olmoe", "qwen36", "deepseek_v4", + "glm", "inkling", "kimi", "olmoe", "qwen36", "deepseek_v4", "minimax_m3", }; for (size_t i = 0; i < sizeof(expected) / sizeof(expected[0]); i++) if (strcmp(expected[i], family_id) == 0) return 1; @@ -13,7 +13,7 @@ static int expected_family(const char *family_id) { } int main(void) { - const size_t required_families = 6; + const size_t required_families = 7; size_t count = coli_segment_conformance_fixture_count(); if (count != required_families) { fprintf(stderr, "segment conformance requires %zu families, found %zu\n", diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 439af24c6..09c622c25 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -252,6 +252,8 @@ def m3_name(name): return None if ".self_attn.index_" in name: return None if name.startswith("language_model."): name = name[len("language_model."):] + name = name.replace(".block_sparse_moe.e_score_correction_bias", + ".mlp.gate.e_score_correction_bias") # the C loader's GLM name name = name.replace(".block_sparse_moe.", ".mlp.") name = name.replace(".w1.weight", ".gate_proj.weight") name = name.replace(".w3.weight", ".up_proj.weight") diff --git a/c/tools/make_m3tiny.py b/c/tools/make_m3tiny.py new file mode 100644 index 000000000..dcaef72f2 --- /dev/null +++ b/c/tools/make_m3tiny.py @@ -0,0 +1,70 @@ +"""Tiny random MiniMax-M3-shaped checkpoint (VL layout, language_model.* prefix) for +engine-vs-oracle validation — the M3 counterpart of glm_tiny. Layer 0 dense + layer 1 +MoE, GQA 4/2 heads with head_dim 8 (deliberately != hidden/heads to exercise the +explicit head_dim path), partial rotary 4 of 8, all 4 experts present, nonzero router +bias (exercises the choice-vs-weight distinction). Convert with: + python3 tools/convert_fp8_to_int4.py --indir m3tiny --outdir m3tiny_i8 --ebits 8 --io-bits 8 +then validate with tools/oracle_m3.py + the engine's REF/TF gate.""" +import json, sys, os +import torch +from safetensors.torch import save_file + +OUT = sys.argv[1] if len(sys.argv) > 1 else "m3tiny" +os.makedirs(OUT, exist_ok=True) +torch.manual_seed(42) + +D, L, H, NK, HD = 16, 2, 4, 2, 8 +ROT, E, TOPK, MI, DI, SI, V = 4, 4, 2, 8, 12, 8, 32 + +t = {} +def w(n, o, i): t[n] = (torch.randn(o, i) * 0.3).to(torch.bfloat16) +def vv(n, o): t[n] = (torch.randn(o) * 0.3).to(torch.bfloat16) + +P = "language_model." +w(P+"model.embed_tokens.weight", V, D) +w(P+"lm_head.weight", V, D) +vv(P+"model.norm.weight", D) +for i in range(L): + Lp = P + f"model.layers.{i}." + vv(Lp+"input_layernorm.weight", D) + vv(Lp+"post_attention_layernorm.weight", D) + w(Lp+"self_attn.q_proj.weight", H*HD, D) + w(Lp+"self_attn.k_proj.weight", NK*HD, D) + w(Lp+"self_attn.v_proj.weight", NK*HD, D) + w(Lp+"self_attn.o_proj.weight", D, H*HD) + vv(Lp+"self_attn.q_norm.weight", HD) + vv(Lp+"self_attn.k_norm.weight", HD) + if i == 0: # dense layer + w(Lp+"mlp.gate_proj.weight", DI, D) + w(Lp+"mlp.up_proj.weight", DI, D) + w(Lp+"mlp.down_proj.weight", D, DI) + else: # MoE layer + M = Lp + "block_sparse_moe." + w(M+"gate.weight", E, D) + t[M+"e_score_correction_bias"] = (torch.randn(E) * 0.1).to(torch.float32) + for e in range(E): + w(M+f"experts.{e}.w1.weight", MI, D) # gate + w(M+f"experts.{e}.w2.weight", D, MI) # down + w(M+f"experts.{e}.w3.weight", MI, D) # up + w(M+"shared_experts.gate_proj.weight", SI, D) + w(M+"shared_experts.up_proj.weight", SI, D) + w(M+"shared_experts.down_proj.weight", D, SI) +w("vision_tower.vision_model.embeddings.patch_embedding.weight", 8, 8) # must be dropped + +save_file(t, os.path.join(OUT, "model-00001-of-00001.safetensors")) +cfg = {"model_type": "minimax_m3_vl", + "text_config": { + "hidden_size": D, "num_hidden_layers": L, "num_attention_heads": H, + "num_key_value_heads": NK, "head_dim": HD, "rotary_dim": ROT, + "partial_rotary_factor": ROT/HD, "intermediate_size": MI, + "dense_intermediate_size": DI, "shared_intermediate_size": SI, + "num_local_experts": E, "num_experts_per_tok": TOPK, "n_shared_experts": 1, + "scoring_func": "sigmoid", "use_routing_bias": True, + "routed_scaling_factor": 2.0, "moe_layer_freq": [0, 1], + "rms_norm_eps": 1e-6, "rope_theta": 5000000, "vocab_size": V, + "use_qk_norm": True, "qk_norm_type": "per_head", "use_gemma_norm": True, + "hidden_act": "swigluoai", "swiglu_alpha": 1.702, "swiglu_limit": 7.0, + "eos_token_id": V-1}, + "vision_config": {"hidden_size": 8}} +json.dump(cfg, open(os.path.join(OUT, "config.json"), "w"), indent=1) +print(f"{OUT}: {len(t)} tensors, D={D} L={L} H={H}/{NK} hd={HD} rot={ROT} E={E} top{TOPK} V={V}") diff --git a/c/tools/oracle_m3.py b/c/tools/oracle_m3.py new file mode 100644 index 000000000..930addceb --- /dev/null +++ b/c/tools/oracle_m3.py @@ -0,0 +1,120 @@ +"""Numpy oracle for the MiniMax-M3 engine port. Reads a CONVERTED tiny container +(so the reference runs on the exact same dequantized weights as the C engine — +differences are implementation bugs, not quantization), teacher-forces a fixed +token sequence with the pinned M3 conventions, and writes the engine's REF gate: + ref_m3.json {prompt_ids, full_ids, tf_pred} + oracle_logits.npy [T,V] +Validate: REF=ref_m3.json TF=1 COLI_MODEL= ./colibri 8 (expect T/T) + +Conventions (transformers modeling_minimax_m3_vl, pinned 2026-07-22): + Gemma RMSNorm x/rms*(1+w) in f32; per-head QK-norm BEFORE partial NEOX RoPE + (first `rotary` dims, split-half within them, inv_freq theta^(-2j/rot)); + GQA repeat_kv h -> h // (H/NK); scale 1/sqrt(head_dim); swigluoai + (gate=min(g,lim), up=clamp(±lim), (up+1)*gate*sigmoid(alpha*gate)); + router sigmoid -> +bias for CHOICE only, raw-sigmoid weights renormalized, + routed_out * routed_scaling + shared_out; pre-norm residual layers.""" +import json, sys, glob +import numpy as np +from safetensors import safe_open + +IND = sys.argv[1] if len(sys.argv) > 1 else "m3tiny_i8" + +T = {} +for p in sorted(glob.glob(f"{IND}/out-*.safetensors")): + with safe_open(p, framework="np") as f: + for k in f.keys(): T[k] = f.get_tensor(k) +cfg = json.load(open(f"{IND}/config.json")) +D, L, H = cfg["hidden_size"], cfg["num_hidden_layers"], cfg["num_attention_heads"] +NK, HD, ROT = cfg["num_key_value_heads"], cfg["head_dim"], cfg["rotary_dim"] +E, TOPK, V = cfg["num_local_experts"], cfg["num_experts_per_tok"], cfg["vocab_size"] +EPS, TH = cfg["rms_norm_eps"], cfg["rope_theta"] +A, LIM, RS = cfg["swiglu_alpha"], cfg["swiglu_limit"], cfg["routed_scaling_factor"] +FD = sum(1 for x in (cfg.get("moe_layer_freq") or []) if x == 0) if cfg.get("moe_layer_freq") else 0 + +def deq(name): + """Dequant a container tensor exactly as the engine reads it (int8 per-row, f32).""" + wq = T[name] + if name + ".qs" not in T: return wq.astype(np.float32) + s = T[name + ".qs"].astype(np.float32) + q = wq.view(np.int8).astype(np.float32) + O = s.shape[0]; q = q.reshape(O, -1) + return q * s[:, None] + +def rms(x, w, eps=EPS): # gemma: (1+w) + r = 1.0 / np.sqrt((x.astype(np.float64)**2).mean(-1, keepdims=True) + eps) + return (x * r * (1.0 + w)).astype(np.float32) + +def rope(vh, pos): # partial split-half NEOX on first ROT dims + h2 = ROT // 2 + j = np.arange(h2, dtype=np.float32) + fr = TH ** (-2.0 * j / ROT) + ang = pos * fr; cs, sn = np.cos(ang), np.sin(ang) + x1, x2 = vh[:h2].copy(), vh[h2:ROT].copy() + vh[:h2] = x1 * cs - x2 * sn + vh[h2:ROT] = x2 * cs + x1 * sn + return vh + +def act(g, u): # swigluoai + g = np.minimum(g, LIM) + u = np.clip(u, -LIM, LIM) + return (u + 1.0) * (g / (1.0 + np.exp(-A * g))) + +def sigmoid(x): return 1.0 / (1.0 + np.exp(-x)) + +ids = [(7 * i + 3) % V for i in range(24)] +S = len(ids) +emb = deq("model.embed_tokens.weight") +x = emb[ids].astype(np.float32) # [S,D] + +for li in range(L): + P = f"model.layers.{li}." + nrm = rms(x, T[P + "input_layernorm.weight"].astype(np.float32)) + q = nrm @ deq(P + "self_attn.q_proj.weight").T # [S,H*HD] + k = nrm @ deq(P + "self_attn.k_proj.weight").T # [S,NK*HD] + v = nrm @ deq(P + "self_attn.v_proj.weight").T + qn = T[P + "self_attn.q_norm.weight"].astype(np.float32) + kn = T[P + "self_attn.k_norm.weight"].astype(np.float32) + q = q.reshape(S, H, HD); k = k.reshape(S, NK, HD); v = v.reshape(S, NK, HD) + for s in range(S): + for h in range(H): q[s, h] = rope(rms(q[s, h], qn), s) + for h in range(NK): k[s, h] = rope(rms(k[s, h], kn), s) + G = H // NK + ctx = np.zeros((S, H, HD), np.float32) + for s in range(S): + for h in range(H): + g = h // G + sc = (q[s, h] @ k[:s+1, g].T) / np.sqrt(HD) + sc = sc - sc.max(); e = np.exp(sc); a = e / e.sum() + ctx[s, h] = a @ v[:s+1, g] + x = x + ctx.reshape(S, H * HD) @ deq(P + "self_attn.o_proj.weight").T + nrm = rms(x, T[P + "post_attention_layernorm.weight"].astype(np.float32)) + if li < FD: # dense + g = nrm @ deq(P + "mlp.gate_proj.weight").T + u = nrm @ deq(P + "mlp.up_proj.weight").T + x = x + act(g, u) @ deq(P + "mlp.down_proj.weight").T + else: # MoE + shared + sg = nrm @ deq(P + "mlp.shared_experts.gate_proj.weight").T + su = nrm @ deq(P + "mlp.shared_experts.up_proj.weight").T + sh = act(sg, su) @ deq(P + "mlp.shared_experts.down_proj.weight").T + logits = nrm @ T[P + "mlp.gate.weight"].astype(np.float32).T + rw = sigmoid(logits.astype(np.float64)).astype(np.float32) + bias = T[P + "mlp.gate.e_score_correction_bias"].astype(np.float32) + routed = np.zeros_like(nrm) + for s in range(S): + choice = rw[s] + bias + top = np.argsort(-choice)[:TOPK] + wts = rw[s][top]; wts = wts / wts.sum() + for t_i, e_i in enumerate(top): + EP = P + f"mlp.experts.{e_i}." + g = nrm[s] @ deq(EP + "gate_proj.weight").T + u = nrm[s] @ deq(EP + "up_proj.weight").T + routed[s] += wts[t_i] * (act(g, u) @ deq(EP + "down_proj.weight").T) + x = x + routed * RS + sh + +fn = T["model.norm.weight"].astype(np.float32) +lo = rms(x, fn) @ deq("lm_head.weight").T # [S,V] +pred = lo.argmax(-1).tolist() +json.dump({"prompt_ids": ids[:4], "full_ids": ids, "tf_pred": pred}, + open("ref_m3.json", "w")) +np.save("oracle_logits.npy", lo) +print(f"oracle: {S} positions, logits [{S},{V}] -> ref_m3.json + oracle_logits.npy") +print("tf_pred:", pred) From 222d8b82763cc8cf49efcf945068c77bdc0e753d Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Wed, 22 Jul 2026 12:32:28 +0200 Subject: [PATCH 03/22] oracle gate: TF_LOGITS full-row dump + TF_DECODE incremental check MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit TF_LOGITS=1 prints every teacher-forcing logit row (oracle bisection); TF_DECODE=1 re-runs the same REF continuation through the incremental S=1 decode path (fresh kv_alloc, prefill np then per-token steps) and compares argmax against tf_pred — validates KV append across calls, which the batch check alone cannot see. M3 tiny results: prefill 24/24 (max logit diff 3.2e-6 vs the numpy oracle with IDOT=0 exact kernels; the idot default's activation quant explains the earlier 23/24), decode 20/20. --- c/colibri.c | 20 ++++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/c/colibri.c b/c/colibri.c index 4744922f5..43f3dd5c2 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -7852,9 +7852,13 @@ static void forward_all(Model *m, const int *ids, int S, int *pred, const int *r layers_forward(m,x,S,0); float *lo=falloc(c->vocab); float *row=falloc(D); + int dump_lo = getenv("TF_LOGITS")!=NULL; /* oracle bisection: full logit rows */ for(int s=0;sfinal_norm, D, c->eps); /* heap row (#183) */ matmul_qt(lo, row, &m->lm_head, 1); + if(dump_lo){ fprintf(stderr,"TFLOGITS %d:",s); + for(int i=0;ivocab;i++) fprintf(stderr," %.6f",lo[i]); + fprintf(stderr,"\n"); } int best=0; float bv=lo[0]; for(int i=1;ivocab;i++) if(lo[i]>bv){bv=lo[i];best=i;} pred[s]=best; if(dbg && pred[s]!=ref[s]) dump_top5_logits(s, lo, c->vocab, ref[s], pred[s]); @@ -11496,6 +11500,22 @@ int main(int argc, char **argv){ if(ok=np) { lg=step(&m,full+i,1,i); } + int best=0; for(int t=1;t Date: Thu, 23 Jul 2026 13:41:04 +0200 Subject: [PATCH 04/22] =?UTF-8?q?vulkan:=20swigluoai=20activation=20in=20t?= =?UTF-8?q?he=20fused=20gate+up=20shader=20=E2=80=94=20enable=20the=20M3?= =?UTF-8?q?=20expert=20tier?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The fused gate+up shader hardcoded silu(gate)*up, so the VK expert tier was gated off for MiniMax-M3 (swigluoai) and its experts ran on the CPU. Add a push-const- selected activation: act=1 computes (up+1)*gate*sigmoid(alpha*gate) with gate/up clamped to +/-limit (alpha 1.702, limit 7.0), matching act_glu and the numpy oracle; act=0 keeps GLM's silu bit-identical. The activation is model-global, set once via coli_vk_set_activation from vk_registry_fill (g_act_swigluoai -> the mode), read into every gate_up dispatch's push constants (both the G and dev2 G2 paths). Un-gates vk_registry_fill for M3. Validated on the real 225 GB int4-g64 container, RX 9070: expert tier fills (384 resident, 10.9 GB, vk 40% of lookups), output stays correct ('Paris' greedy + coherent prose), decode 0.64 -> 0.91 tok/s (+42%) with the tier active. Attention (GQA) and the shared expert still run on the CPU (P4). --- c/backend_vulkan.c | 17 ++++++++++++++++- c/backend_vulkan.h | 2 ++ c/colibri.c | 9 ++++----- c/shaders/qmatmul_gate_up.comp | 15 ++++++++++++++- 4 files changed, 36 insertions(+), 7 deletions(-) diff --git a/c/backend_vulkan.c b/c/backend_vulkan.c index d779607d8..87c39e8bc 100644 --- a/c/backend_vulkan.c +++ b/c/backend_vulkan.c @@ -104,7 +104,19 @@ static struct { float prio; /* priority applied to the NEXT allocations (class knob) */ } G; -struct PC { int fmt, S, I, O, rowWords, gs; }; +struct PC { int fmt, S, I, O, rowWords, gs; int act; float alpha, limit; }; + +/* Fused gate+up activation, model-global (all experts + the shared/dense MLP share it): + * 0 = silu(gate)*up (GLM), 1 = swigluoai (MiniMax-M3). Set once via coli_vk_set_activation; + * applied in every gate_up dispatch's push constants. File-static so both device contexts + * (G and the dev2 G2) read the same value. */ +static int vk_act; +static float vk_alpha = 1.702f, vk_limit = 7.0f; +void coli_vk_set_activation(int act, float alpha, float limit) { + vk_act = act; + if (alpha > 0) vk_alpha = alpha; + if (limit > 0) vk_limit = limit; +} struct PCN { int S, D; float eps; }; /* Push constants of the absorb attention kernel (must match attention_absorb.comp). */ struct PCAttn { int fmt, S, H, Q, R, V, K, st0, T, rowWords, cap; float scale; int gs; }; @@ -685,6 +697,7 @@ int coli_vk_gate_up(ColiVkTensor **gate, ColiVkTensor **up, float *hidden, const vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gu, 0, 1, &G.dset_gu, 0, NULL); struct PC pc = {fmt, S, D, I, tg->rowWords, tg->gs}; // PC.I = input D, PC.O = moe_inter I + pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdPushConstants(G.cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G.cmd, (uint32_t)((I + 7) / 8), (uint32_t)S, 1); VKCHECK(vkEndCommandBuffer(G.cmd), "endCmd"); @@ -777,6 +790,7 @@ static int eg_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *up vkCmdBindPipeline(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); for (int c = 0; c < count; c++) { struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs}; + pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdBindDescriptorSets(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gu, 0, 1, &G.eg_gu[c], 0, NULL); vkCmdPushConstants(G.eg_cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G.eg_cmd, (uint32_t)((I + 7) / 8), (uint32_t)rows[c], 1); @@ -1119,6 +1133,7 @@ static int eg2_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *u vkCmdBindPipeline(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.pipe_gu); for (int c = 0; c < count; c++) { struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs}; + pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdBindDescriptorSets(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.plyt_gu, 0, 1, &G2.gu[c], 0, NULL); vkCmdPushConstants(G2.cmd, G2.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G2.cmd, (uint32_t)((I + 7) / 8), (uint32_t)rows[c], 1); diff --git a/c/backend_vulkan.h b/c/backend_vulkan.h index 30b617d67..673de3bd2 100644 --- a/c/backend_vulkan.h +++ b/c/backend_vulkan.h @@ -29,6 +29,8 @@ void coli_vk_mem_info(size_t *used_bytes, size_t *tensor_count); * mem_budget reports device-local usage/budget in GB (VK_EXT_memory_budget); * returns 0 if unavailable. */ void coli_vk_alloc_priority(float p); +/* Fused gate+up activation (model-global): 0 = silu(gate)*up (GLM), 1 = swigluoai (M3). */ +void coli_vk_set_activation(int act, float alpha, float limit); int coli_vk_mem_budget(double *used_gb, double *budget_gb); /* y[S,O] = (x[S,I] @ dequant(W[O,I])^T) * scale[O]. diff --git a/c/colibri.c b/c/colibri.c index 43f3dd5c2..22334479e 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -9432,11 +9432,10 @@ static void vk_dense_preload(Model *m){ static void vk_registry_fill(Model *m){ Cfg *c=&m->c; int E=c->n_experts, NL=c->n_layers; if(!g_vulkan || g_vk_budget<=0) return; - if(c->arch!=ARCH_GLM){ /* expert-group shaders hardcode silu */ - fprintf(stderr,"[VK] expert tier disabled for this architecture (shader activation " - "is silu; swigluoai shader variant pending) — experts run on the CPU\n"); - return; - } + /* The fused gate+up shader selects its activation via a push constant — silu for GLM, + * swigluoai for MiniMax-M3 — so the expert tier runs on any of our arches. Set the + * model's choice once here, before the tier fills (all gate_up dispatches read it). */ + coli_vk_set_activation(g_act_swigluoai, c->swiglu_alpha, c->swiglu_limit); int64_t nz=0; for(int i=0;ieusage[i]) for(int e=0;eeusage[i][e]) nz++; if(!nz){ fprintf(stderr,"[VK] expert tier: no usage history yet — tier empty this run " diff --git a/c/shaders/qmatmul_gate_up.comp b/c/shaders/qmatmul_gate_up.comp index 8db58d047..56fe2d4a0 100644 --- a/c/shaders/qmatmul_gate_up.comp +++ b/c/shaders/qmatmul_gate_up.comp @@ -25,6 +25,9 @@ layout(push_constant) uniform P { int S, I, O; // I = input dim (hidden), O = moe_inter int rowWords; // uint32 words per weight row (of the I-dim rows) int gs; // fmt=4 group size (multiple of 8; host gates otherwise) + int act; // gate activation: 0 = silu(gate)*up (GLM), 1 = swigluoai (MiniMax-M3) + float alpha; // swigluoai: sigmoid(alpha*gate), alpha = 1.702 + float limit; // swigluoai: gate=min(g,limit), up=clamp(u,+/-limit), limit = 7.0 } p; shared float xsh[6144]; @@ -99,6 +102,16 @@ void main() { } float gt = subgroupAdd(g), ut = subgroupAdd(u); if (p.fmt != 5 && p.fmt != 4) { gt *= gscale[o]; ut *= uscale[o]; } // per-row formats only - if (lane == 0) hid[s * p.O + o] = (gt / (1.0 + exp(-gt))) * ut; // silu(gate)*up + if (lane == 0) { + float h; + if (p.act == 1) { // swigluoai (MiniMax-M3): (up+1)*gate*sigmoid(alpha*gate) + float gc = min(gt, p.limit); // gate clamped above + float uc = clamp(ut, -p.limit, p.limit); // up clamped both sides + h = (uc + 1.0) * gc / (1.0 + exp(-p.alpha * gc)); + } else { // silu(gate)*up (GLM) + h = (gt / (1.0 + exp(-gt))) * ut; + } + hid[s * p.O + o] = h; + } } } From ca734b253da7c7bc6e36bf0bdbfb8a02043ee809 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 23 Jul 2026 14:40:48 +0200 Subject: [PATCH 05/22] vulkan: GQA attention core for MiniMax-M3 (COLI_VK_ATTN) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit New attention_gqa.comp: one workgroup per (query row, head) computes scores/softmax/weighted-V for GQA (query head h reads KV group h/(H/NK)), reading the persistent on-device K/V mirror — K in the L buffer, V in the R buffer, both NK*hd wide (the mirror already stores two arbitrary-width rows; the absorb shader's R<=64 cap is its own, not the mirror's). coli_vk_gqa_attn mirrors the absorb dispatch (5 bindings, one submit+fence+readback per layer). Wired into attention_gqa behind COLI_VK_ATTN, single-sequence decode only, CPU fallback on any failure. Validated on the real 225 GB container, RX 9070: output correct ('Paris' greedy + coherent prose) with the core active. HONEST PERF NOTE: it is throughput-NEUTRAL on this box and stays opt-in/default-off, because the core is only ~0.7s of an ~8.6s decode-attention cost — the projections dominate (q/k/v proj+norm+rope ~4.8s, o-proj ~2.9s) and are still on the CPU. The per- layer submit/readback slightly exceeds the tiny core it replaces. Offloading the projections (fused, GLM-qprep style, to amortize submits) is the actual lever and the next step. -Wno-missing-field-initializers quiets the shared struct PC's zero-filled tail. --- c/Makefile | 12 +++--- c/backend_vulkan.c | 68 ++++++++++++++++++++++++++++++- c/backend_vulkan.h | 4 ++ c/colibri.c | 23 +++++++++++ c/shaders/attention_gqa.comp | 78 ++++++++++++++++++++++++++++++++++++ 5 files changed, 177 insertions(+), 8 deletions(-) create mode 100644 c/shaders/attention_gqa.comp diff --git a/c/Makefile b/c/Makefile index 290b75dcd..06fd6882c 100644 --- a/c/Makefile +++ b/c/Makefile @@ -62,7 +62,7 @@ $(warning libomp not found: building single-threaded. For multithreading: brew i OMPC = OMPL = endif -CFLAGS = -O3 $(OMPC) -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) +CFLAGS = -O3 $(OMPC) -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) # Opt-in: ARCH=native appends -mcpu=native (arm64 clang uses -mcpu, not -march), # which unlocks the i8mm SMMLA int8/int4 dot kernels in colibri.c. ARCH unset -> # no -mcpu, default build byte-identical. Apple clang knows apple-m4 / native. @@ -96,7 +96,7 @@ else ifneq ($(IS_WIN),) # a v3 build simply compiles out the VNNI path - safe on any x86-64. CC = gcc ARCH ?= x86-64-v3 -CFLAGS = -D_FILE_OFFSET_BITS=64 -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) +CFLAGS = -D_FILE_OFFSET_BITS=64 -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) # -lpsapi: compat.h calls GetProcessMemoryInfo (rss_gb). It's linked via # #pragma comment(lib,"psapi.lib") for MSVC, but MinGW gcc ignores that pragma # (warns), so link psapi explicitly or the build fails undefined-reference on @@ -115,7 +115,7 @@ ifneq (,$(PPC64)) # (validated token-exact vs the transformers oracle on a POWER8 S824). CC = gcc ARCH ?= native -CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) +CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto @@ -168,7 +168,7 @@ ARCHFLAG := -march=$(ARCH) else ARCHFLAG := -mcpu=$(ARCH) endif -CFLAGS = -O3 $(ARCHFLAG) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) +CFLAGS = -O3 $(ARCHFLAG) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto @@ -183,7 +183,7 @@ CC = gcc # ARCH=x86-64 -> massima compatibilita' (niente AVX2: usa il path scalare di fallback). # -pthread: Linux lo tira dentro via -fopenmp, i *BSD no e le pthread_* non risolvono (#219). ARCH ?= native -CFLAGS = -O3 -march=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) +CFLAGS = -O3 -march=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto @@ -602,7 +602,7 @@ ifeq ($(VK),1) CFLAGS += -DCOLI_VULKAN LDFLAGS += -lvulkan VK_OBJ = backend_vulkan.o -VK_SPV = shaders/qmatmul.spv shaders/qmatmul_gate_up.spv shaders/attention_absorb.spv shaders/rmsnorm.spv +VK_SPV = shaders/qmatmul.spv shaders/qmatmul_gate_up.spv shaders/attention_absorb.spv shaders/attention_gqa.spv shaders/rmsnorm.spv endif all: colibri$(EXE) diff --git a/c/backend_vulkan.c b/c/backend_vulkan.c index 87c39e8bc..58637367e 100644 --- a/c/backend_vulkan.c +++ b/c/backend_vulkan.c @@ -61,6 +61,8 @@ static struct { /* MLA absorb attention core (7 bindings): q, W, scales, Lcache, Rcache, scores, ctx */ VkShaderModule shader_att; VkDescriptorSetLayout dsl_att; VkPipelineLayout plyt_att; VkPipeline pipe_att; VkDescriptorPool dpool_att; VkDescriptorSet dset_att; + VkShaderModule shader_gqa; VkDescriptorSetLayout dsl_gqa; VkPipelineLayout plyt_gqa; + VkPipeline pipe_gqa; VkDescriptorPool dpool_gqa; VkDescriptorSet dset_gqa; VkCommandPool cpool; VkCommandBuffer cmd; VkFence fence; @@ -120,6 +122,7 @@ void coli_vk_set_activation(int act, float alpha, float limit) { struct PCN { int S, D; float eps; }; /* Push constants of the absorb attention kernel (must match attention_absorb.comp). */ struct PCAttn { int fmt, S, H, Q, R, V, K, st0, T, rowWords, cap; float scale; int gs; }; +struct PCGqa { int S, H, NK, hd, st0, T, cap; float scale; }; /* M3 GQA attention core */ static int pick_memtype(VkPhysicalDevice phys) { VkPhysicalDeviceMemoryProperties m; @@ -439,6 +442,12 @@ int coli_vk_init(const char *spv_path) { if (G.shader_att && !build_pipeline(G.dev, 7, sizeof(struct PCAttn), G.shader_att, &G.dsl_att, &G.plyt_att, &G.pipe_att, &G.dpool_att, &G.dset_att)) return 0; + /* Optional GQA attention core (MiniMax-M3): Q, K-cache, V-cache, scores, ctx. */ + char gqa_path[512]; derive_dir_file(spv_path, "attention_gqa.spv", gqa_path, sizeof(gqa_path)); + G.shader_gqa = load_spv(G.dev, gqa_path); + if (G.shader_gqa && !build_pipeline(G.dev, 5, sizeof(struct PCGqa), G.shader_gqa, &G.dsl_gqa, &G.plyt_gqa, &G.pipe_gqa, &G.dpool_gqa, &G.dset_gqa)) + return 0; + VkCommandPoolCreateInfo cpci = {.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO, .flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT, .queueFamilyIndex = G.qfam}; VKCHECK(vkCreateCommandPool(G.dev, &cpci, NULL, &G.cpool), "cmdPool"); @@ -452,8 +461,9 @@ int coli_vk_init(const char *spv_path) { G.ready = 1; VkPhysicalDeviceProperties p; vkGetPhysicalDeviceProperties(G.phys, &p); - fprintf(stderr, "[VK] ready: %s, compute qfam %u, memtype %u%s%s\n", p.deviceName, G.qfam, G.memtype, - G.shader_gu ? ", fused gate+up" : "", G.shader_att ? ", absorb attention" : ""); + fprintf(stderr, "[VK] ready: %s, compute qfam %u, memtype %u%s%s%s\n", p.deviceName, G.qfam, G.memtype, + G.shader_gu ? ", fused gate+up" : "", G.shader_att ? ", absorb attention" : "", + G.shader_gqa ? ", GQA attention" : ""); return 1; } @@ -1285,6 +1295,53 @@ int coli_vk_attention_absorb(ColiVkTensor **kvb, const void *w, const float *sc, return 1; } +/* Decode GQA attention core for S causal query rows (MiniMax-M3), one submit: + * ctx[s,h,:] = softmax((q[s,h] . K_t[g]) * scale) weighted V_t[g], g = h/(H/NK). + * K and V rows [st0, T) must already be mirrored via coli_vk_kv_row (K in the L + * buffer, V in the R buffer, both NK*hd wide). Queries and cache rows arrive + * already normed+roped. Returns 0 -> caller falls back to the CPU core. */ +int coli_vk_gqa_attn(float *ctx, const float *q, int layer, int S, int H, int NK, int hd, + int st0, int T, float scale) { + if (!G.ready || !G.pipe_gqa || S < 1 || H < 1 || NK < 1 || layer < 0 || layer >= VK_KV_LAYERS) return 0; + if (hd > 256 || H % NK != 0 || st0 < 0 || T - S - st0 < 0) return 0; /* qsh[256] limit */ + int KVd = NK * hd; + VkKvLayer *kv = &G.kv[layer]; + if (!kv->bl || kv->rows < T || kv->K != KVd || kv->R != KVd) return 0; + int cap = T - st0; + size_t qb = (size_t)S * H * hd * 4, cb = (size_t)S * H * hd * 4; + size_t sb = (size_t)S * H * cap * 4; + if (!scratch_reserve(&G.x, qb) || !scratch_reserve_mt(&G.y, cb, G.memtype_cached) || + !scratch_reserve(&G.att_sc, sb)) return 0; /* y (ctx) is read back -> cached */ + memcpy(G.x.ptr, q, qb); + + VkDescriptorBufferInfo bi[5] = { + {.buffer = G.x.buf, .range = VK_WHOLE_SIZE}, {.buffer = kv->bl, .range = VK_WHOLE_SIZE}, + {.buffer = kv->br, .range = VK_WHOLE_SIZE}, {.buffer = G.att_sc.buf, .range = VK_WHOLE_SIZE}, + {.buffer = G.y.buf, .range = VK_WHOLE_SIZE}}; + wr_desc(G.dset_gqa, 5, bi); + + VKCHECK(vkResetCommandBuffer(G.cmd, 0), "resetCmd"); + VkCommandBufferBeginInfo begin = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; + VKCHECK(vkBeginCommandBuffer(G.cmd, &begin), "beginCmd"); + vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gqa); + vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gqa, 0, 1, &G.dset_gqa, 0, NULL); + struct PCGqa pc = {S, H, NK, hd, st0, T, cap, scale}; + vkCmdPushConstants(G.cmd, G.plyt_gqa, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); + vkCmdDispatch(G.cmd, (uint32_t)H, (uint32_t)S, 1); /* one workgroup per (head, row) */ + VKCHECK(vkEndCommandBuffer(G.cmd), "endCmd"); + + VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &G.cmd}; + VKCHECK(vkResetFences(G.dev, 1, &G.fence), "resetFence"); + double vp0 = G.eg_prof ? vk_now() : 0; + VKCHECK(vkQueueSubmit(G.queue, 1, &si, G.fence), "queueSubmit"); + if (G.eg_prof) { double vp1 = vk_now(); g_vsub_ms += vp1 - vp0; vp0 = vp1; } + if (vk_fence_wait(G.dev, G.fence) != VK_SUCCESS) { G.ready = 0; return 0; } + if (G.eg_prof) { g_vwait_ms += vk_now() - vp0; vkprof_tick(); } + memcpy(ctx, G.y.ptr, cb); + G.cmd_ready = 0; G.bound_tensor = NULL; /* the shared command buffer/binding was clobbered */ + return 1; +} + /* Two resident matmuls sharing the SAME input x in ONE submit (q_a + kv_a in the * attention prologue): one x staging, two dispatches, one fence — replaces two * full submit+wait roundtrips. Outputs y1 [S,O1] and y2 [S,O2] read back from @@ -1585,6 +1642,13 @@ void coli_vk_shutdown(void) { vkDestroyDescriptorSetLayout(G.dev, G.dsl_att, NULL); vkDestroyShaderModule(G.dev, G.shader_att, NULL); } + if (G.shader_gqa) { + vkDestroyDescriptorPool(G.dev, G.dpool_gqa, NULL); + vkDestroyPipeline(G.dev, G.pipe_gqa, NULL); + vkDestroyPipelineLayout(G.dev, G.plyt_gqa, NULL); + vkDestroyDescriptorSetLayout(G.dev, G.dsl_gqa, NULL); + vkDestroyShaderModule(G.dev, G.shader_gqa, NULL); + } for (VkWArena *a = g_warena; a;) { /* weight arenas: unmapped/freed with the device */ VkWArena *nx = a->next; vkUnmapMemory(G.dev, a->mem); vkFreeMemory(G.dev, a->mem, NULL); diff --git a/c/backend_vulkan.h b/c/backend_vulkan.h index 673de3bd2..3a12f3378 100644 --- a/c/backend_vulkan.h +++ b/c/backend_vulkan.h @@ -31,6 +31,10 @@ void coli_vk_mem_info(size_t *used_bytes, size_t *tensor_count); void coli_vk_alloc_priority(float p); /* Fused gate+up activation (model-global): 0 = silu(gate)*up (GLM), 1 = swigluoai (M3). */ void coli_vk_set_activation(int act, float alpha, float limit); +/* Decode GQA attention core (MiniMax-M3): K/V rows mirrored via coli_vk_kv_row + * (K in the L buffer, V in the R buffer, both NK*hd wide). Returns 0 -> CPU fallback. */ +int coli_vk_gqa_attn(float *ctx, const float *q, int layer, int S, int H, int NK, int hd, + int st0, int T, float scale); int coli_vk_mem_budget(double *used_gb, double *budget_gb); /* y[S,O] = (x[S,I] @ dequant(W[O,I])^T) * scale[O]. diff --git a/c/colibri.c b/c/colibri.c index 22334479e..d83a85311 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -4229,6 +4229,29 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po memcpy(ks->Rc[layer]+(int64_t)pos*KVd, vsr, (size_t)KVd*sizeof(float)); } double tc0=now_s(); + int vk_core=0; (void)vk_core; +#ifdef COLI_VULKAN + /* Vulkan GQA core (COLI_VK_ATTN=1): scores/softmax/weighted-V for all S x H in + * ONE submit per layer, reading the persistent on-device K/V mirror (K in the L + * buffer, V in the R buffer; rows appended incrementally, vk_kv_valid watermark, + * invalidated like the CUDA/MLA shadow on rewrite/rebind/resize). Single-sequence + * decode only (no ragged mux); falls back to the CPU core on any failure. */ + if(g_vk_attn && !kvs && !positions && S<=4 && layern_layers && + m->vk_kv_valid && m->kv->Lc[layer] && m->kv->Rc[layer]){ + int st0=m->kv_start[layer], T=pos_base+S; + if(T<=m->max_t && coli_vk_kv_ensure(layer,m->max_t,KVd,KVd)){ + int ok=1; + for(int t=m->vk_kv_valid[layer];tkv->Lc[layer]+(int64_t)t*KVd, + m->kv->Rc[layer]+(int64_t)t*KVd); + if(ok){ + m->vk_kv_valid[layer]=T; + if(coli_vk_gqa_attn(ctx,q,layer,S,H,NK,hd,st0,T,c->attn_scale)) vk_core=1; + } + } + } +#endif + if(!vk_core) for(int s=0;skv; int pos=positions?positions[s]:pos_base+s; diff --git a/c/shaders/attention_gqa.comp b/c/shaders/attention_gqa.comp new file mode 100644 index 000000000..df6057d79 --- /dev/null +++ b/c/shaders/attention_gqa.comp @@ -0,0 +1,78 @@ +#version 450 +// GQA attention core (decode), the VK equivalent of colibri's attention_gqa CPU +// loop (MiniMax-M3). One workgroup per (query row s = WorkGroupID.y, head +// h = WorkGroupID.x); the whole core for one head runs in a single dispatch: +// g = h / (H/NK) (repeat_kv group) +// score_t = (q[s,h] . K_t[g]) * scale t in [st0, T-S+s] (causal) +// p = softmax(score) +// ctx[i] = sum_t p_t * V_t[g][i] +// K and V are the PERSISTENT device-side caches, each row KVd = NK*hd wide, rows +// appended by the host each token (absolute position indexing, kv_start windows +// skip rows). Queries and cache rows arrive already normed + roped (host writes +// them that way) and are plain f32 — no dequant, no weight absorption. +#extension GL_KHR_shader_subgroup_arithmetic : require +layout(local_size_x = 256) in; + +layout(std430, binding = 0) readonly buffer QB { float q[]; }; // [S,H,hd] roped +layout(std430, binding = 1) readonly buffer KB { float kc[]; }; // K cache [max_t, NK*hd] +layout(std430, binding = 2) readonly buffer VB { float vc[]; }; // V cache [max_t, NK*hd] +layout(std430, binding = 3) buffer SC { float sc[]; }; // scores [S,H,cap] scratch +layout(std430, binding = 4) writeonly buffer CB { float ctx[]; }; // [S,H,hd] out + +layout(push_constant) uniform P { + int S, H, NK, hd; // heads, kv-heads, head_dim + int st0, T; // attend rows [st0, T-S+s]; T = absolute rows incl. final query + int cap; // score scratch stride per (s,h) = T-st0 + float scale; // 1/sqrt(head_dim) +} p; + +shared float qsh[256]; // this head's query row +shared float red[256]; // reduction scratch + +void main() { + int h = int(gl_WorkGroupID.x); + int s = int(gl_WorkGroupID.y); + int tid = int(gl_LocalInvocationID.x); + int nth = int(gl_WorkGroupSize.x); + int lane = int(gl_SubgroupInvocationID); + int sgsize = int(gl_SubgroupSize); + int sg = int(gl_SubgroupID); + int nsg = nth / sgsize; + + int G = p.H / p.NK; + int gof = (h / G) * p.hd; // this query head's KV group slice in a row + int KVd = p.NK * p.hd; + int nt = (p.T - p.S + s + 1) - p.st0; // causal window for query s + int qoff = (s * p.H + h) * p.hd; + int scoff = (s * p.H + h) * p.cap; + + for (int i = tid; i < p.hd; i += nth) qsh[i] = q[qoff + i]; + barrier(); + + /* 1) scores: one subgroup per cached token, lanes split the hd dot */ + float lmax = -3.0e38; + for (int j = sg; j < nt; j += nsg) { + int t = p.st0 + j; + float a = 0.0; + for (int i = lane; i < p.hd; i += sgsize) a += qsh[i] * kc[t * KVd + gof + i]; + float tot = subgroupAdd(a) * p.scale; + if (lane == 0) { sc[scoff + j] = tot; lmax = max(lmax, tot); } + } + red[tid] = lmax; barrier(); + for (int o = nth / 2; o > 0; o >>= 1) { if (tid < o) red[tid] = max(red[tid], red[tid + o]); barrier(); } + float wmax = red[0]; barrier(); + + /* 2) softmax: exp in place + workgroup sum */ + float lsum = 0.0; + for (int j = tid; j < nt; j += nth) { float e = exp(sc[scoff + j] - wmax); sc[scoff + j] = e; lsum += e; } + red[tid] = lsum; barrier(); + for (int o = nth / 2; o > 0; o >>= 1) { if (tid < o) red[tid] += red[tid + o]; barrier(); } + float inv = 1.0 / red[0]; barrier(); + + /* 3) weighted value: ctx[i] = sum_t p_t V_t[g][i] — thread-parallel over head dim */ + for (int i = tid; i < p.hd; i += nth) { + float acc = 0.0; + for (int j = 0; j < nt; j++) acc += sc[scoff + j] * vc[(p.st0 + j) * KVd + gof + i]; + ctx[(s * p.H + h) * p.hd + i] = acc * inv; + } +} From deb50eb7a926066185c11ed66776dbedc74e160d Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 23 Jul 2026 15:04:31 +0200 Subject: [PATCH 06/22] vulkan: fuse the GQA core with the o-projection (one submit, ctx on-device) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit coli_vk_gqa_attn_project runs the GQA core and the o-projection in one command buffer: the core writes ctx into the att_ctx device scratch (never read back), a compute barrier, then the o-proj matmul reads it and only [S,hidden] returns to the host — the absorb_project pattern, for GQA. attention_gqa tries it first (fills out directly, skips the CPU o-proj), then the core-only path, then CPU. Validated correct ('Paris' greedy). Profile confirms the offload: decode output-projection drops to 0.000s on the CPU, and the core+o-proj VK dispatch (~2.3s/64tok) replaces CPU core+o-proj (~3.9s) — a ~1.6s attention saving. HONEST e2e NOTE: still throughput-neutral on this box (fused 1.68/1.71 vs CPU 1.70/1.73). M3 decode here is expert-matmul + disk bound (expert-matmul ~16-17s, disk-wait ~6s, attention ~9s of a ~34s decode; the Zen2 CPU is AVX2-maxed with no VNNI, and the 16 GB VRAM caps the expert tier at 512). The saving is real but masked by expert/disk variance. Kept behind COLI_VK_ATTN (default off for M3); it pays off on a compute-bound or faster-disk box and sets up full projection fusion. --- c/backend_vulkan.c | 63 ++++++++++++++++++++++++++++++++++++++++++++++ c/backend_vulkan.h | 4 +++ c/colibri.c | 15 ++++++++--- 3 files changed, 78 insertions(+), 4 deletions(-) diff --git a/c/backend_vulkan.c b/c/backend_vulkan.c index 58637367e..cd9980789 100644 --- a/c/backend_vulkan.c +++ b/c/backend_vulkan.c @@ -1342,6 +1342,69 @@ int coli_vk_gqa_attn(float *ctx, const float *q, int layer, int S, int H, int NK return 1; } +/* Fused GQA core + o-projection in ONE submit: the core writes ctx into a device + * scratch (att_ctx, never read back), a barrier, then the o-projection matmul reads + * it and only the final output [S,Dout] returns to the host. Offloads the o-proj + * (the single largest attention cost after the projections) with no extra submit + * over the core alone. Returns 0 -> caller runs the CPU o-proj (or CPU core). */ +int coli_vk_gqa_attn_project(float *out, const float *q, ColiVkTensor **ot, const void *ow, + const float *osc, int ofmt, int ogrp, int layer, int S, int H, + int NK, int hd, int st0, int T, float scale, int Dout) { + if (!G.ready || !G.pipe_gqa || !G.pipe || S < 1 || H < 1 || NK < 1 || layer < 0 || layer >= VK_KV_LAYERS) return 0; + if (hd > 256 || H % NK != 0 || st0 < 0 || T - S - st0 < 0 || Dout < 1) return 0; + int KVd = NK * hd; + VkKvLayer *kv = &G.kv[layer]; + if (!kv->bl || kv->rows < T || kv->K != KVd || kv->R != KVd) return 0; + if (!upload_tensor(ot, ow, osc, ofmt, H * hd, Dout, ogrp)) return 0; /* o: [Dout, H*hd] */ + ColiVkTensor *to = *ot; + int cap = T - st0; + size_t qb = (size_t)S * H * hd * 4, cb = (size_t)S * H * hd * 4; + size_t sb = (size_t)S * H * cap * 4, ob = (size_t)S * Dout * 4; + if (!scratch_reserve(&G.x, qb) || !scratch_reserve(&G.att_ctx, cb) || + !scratch_reserve(&G.att_sc, sb) || !scratch_reserve_mt(&G.y, ob, G.memtype_cached)) return 0; + memcpy(G.x.ptr, q, qb); + + VkDescriptorBufferInfo bi[5] = { /* core: Q, K, V, scores, ctx(->att_ctx device) */ + {.buffer = G.x.buf, .range = VK_WHOLE_SIZE}, {.buffer = kv->bl, .range = VK_WHOLE_SIZE}, + {.buffer = kv->br, .range = VK_WHOLE_SIZE}, {.buffer = G.att_sc.buf, .range = VK_WHOLE_SIZE}, + {.buffer = G.att_ctx.buf, .range = VK_WHOLE_SIZE}}; + wr_desc(G.dset_gqa, 5, bi); + VkDescriptorBufferInfo oi[4] = { /* o-proj: att_ctx, o_w, o_s, out */ + {.buffer = G.att_ctx.buf, .range = VK_WHOLE_SIZE}, {.buffer = to->wbuf, .range = VK_WHOLE_SIZE}, + {.buffer = to->sbuf, .range = VK_WHOLE_SIZE}, {.buffer = G.y.buf, .range = VK_WHOLE_SIZE}}; + wr_desc(G.dset, 4, oi); + + VKCHECK(vkResetCommandBuffer(G.cmd, 0), "resetCmd"); + VkCommandBufferBeginInfo begin = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; + VKCHECK(vkBeginCommandBuffer(G.cmd, &begin), "beginCmd"); + vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gqa); + vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gqa, 0, 1, &G.dset_gqa, 0, NULL); + struct PCGqa pc = {S, H, NK, hd, st0, T, cap, scale}; + vkCmdPushConstants(G.cmd, G.plyt_gqa, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); + vkCmdDispatch(G.cmd, (uint32_t)H, (uint32_t)S, 1); + VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, + .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; + vkCmdPipelineBarrier(G.cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, NULL, 0, NULL); + vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); + vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); + struct PC opc = {ofmt, S, H * hd, Dout, to->rowWords, to->gs}; + vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(opc), &opc); + vkCmdDispatch(G.cmd, (uint32_t)((Dout + 7) / 8), (uint32_t)S, 1); + VKCHECK(vkEndCommandBuffer(G.cmd), "endCmd"); + + VkSubmitInfo si = {.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO, .commandBufferCount = 1, .pCommandBuffers = &G.cmd}; + VKCHECK(vkResetFences(G.dev, 1, &G.fence), "resetFence"); + double vp0 = G.eg_prof ? vk_now() : 0; + VKCHECK(vkQueueSubmit(G.queue, 1, &si, G.fence), "queueSubmit"); + if (G.eg_prof) { double vp1 = vk_now(); g_vsub_ms += vp1 - vp0; vp0 = vp1; } + if (vk_fence_wait(G.dev, G.fence) != VK_SUCCESS) { G.ready = 0; return 0; } + if (G.eg_prof) { g_vwait_ms += vk_now() - vp0; vkprof_tick(); } + memcpy(out, G.y.ptr, ob); + G.cmd_ready = 0; G.bound_tensor = NULL; /* the shared command buffer/binding was clobbered */ + return 1; +} + /* Two resident matmuls sharing the SAME input x in ONE submit (q_a + kv_a in the * attention prologue): one x staging, two dispatches, one fence — replaces two * full submit+wait roundtrips. Outputs y1 [S,O1] and y2 [S,O2] read back from diff --git a/c/backend_vulkan.h b/c/backend_vulkan.h index 3a12f3378..89980725f 100644 --- a/c/backend_vulkan.h +++ b/c/backend_vulkan.h @@ -35,6 +35,10 @@ void coli_vk_set_activation(int act, float alpha, float limit); * (K in the L buffer, V in the R buffer, both NK*hd wide). Returns 0 -> CPU fallback. */ int coli_vk_gqa_attn(float *ctx, const float *q, int layer, int S, int H, int NK, int hd, int st0, int T, float scale); +/* Fused GQA core + o-projection: ctx stays on-device, only [S,Dout] returns. */ +int coli_vk_gqa_attn_project(float *out, const float *q, ColiVkTensor **ot, const void *ow, + const float *osc, int ofmt, int ogrp, int layer, int S, int H, + int NK, int hd, int st0, int T, float scale, int Dout); int coli_vk_mem_budget(double *used_gb, double *budget_gb); /* y[S,O] = (x[S,I] @ dequant(W[O,I])^T) * scale[O]. diff --git a/c/colibri.c b/c/colibri.c index d83a85311..6e88b6167 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -4229,13 +4229,14 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po memcpy(ks->Rc[layer]+(int64_t)pos*KVd, vsr, (size_t)KVd*sizeof(float)); } double tc0=now_s(); - int vk_core=0; (void)vk_core; + int vk_core=0, vk_projected=0; (void)vk_core; (void)vk_projected; #ifdef COLI_VULKAN /* Vulkan GQA core (COLI_VK_ATTN=1): scores/softmax/weighted-V for all S x H in * ONE submit per layer, reading the persistent on-device K/V mirror (K in the L * buffer, V in the R buffer; rows appended incrementally, vk_kv_valid watermark, * invalidated like the CUDA/MLA shadow on rewrite/rebind/resize). Single-sequence - * decode only (no ragged mux); falls back to the CPU core on any failure. */ + * decode only (no ragged mux); falls back to the CPU core on any failure. The + * fused variant keeps ctx on-device and runs the o-projection in the same submit. */ if(g_vk_attn && !kvs && !positions && S<=4 && layern_layers && m->vk_kv_valid && m->kv->Lc[layer] && m->kv->Rc[layer]){ int st0=m->kv_start[layer], T=pos_base+S; @@ -4246,7 +4247,13 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po m->kv->Rc[layer]+(int64_t)t*KVd); if(ok){ m->vk_kv_valid[layer]=T; - if(coli_vk_gqa_attn(ctx,q,layer,S,H,NK,hd,st0,T,c->attn_scale)) vk_core=1; + if(VK_FMT_OK(&l->o)&& + coli_vk_gqa_attn_project(out,q,&l->o.vk, + l->o.fmt==1?(const void*)l->o.q8:(const void*)l->o.q4, + l->o.s,l->o.fmt,l->o.gs,layer,S,H,NK,hd,st0,T,c->attn_scale,c->hidden)) + vk_core=vk_projected=1; + else if(coli_vk_gqa_attn(ctx,q,layer,S,H,NK,hd,st0,T,c->attn_scale)) + vk_core=1; } } } @@ -4285,7 +4292,7 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po } m->t_acore+=now_s()-tc0; double to0=now_s(); - matmul_qt(out,ctx,&l->o,S); + if(!vk_projected) matmul_qt(out,ctx,&l->o,S); /* VK fused path already filled out */ m->t_aout+=now_s()-to0; free(q); free(k); free(v); free(ctx); m->t_attn += now_s()-ta0; From b0e9e5c715fee234f8ea61d34ef785cde4295556 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 23 Jul 2026 16:28:15 +0200 Subject: [PATCH 07/22] engine: MiniMax-M3 MSA (Lightning Indexer block-sparse attention) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit M3's native attention is block-sparse beyond 2048 tokens: a small 4-head indexer scores every key, max-pools into 128-token blocks, and the main GQA attention only reads the top-16 blocks (+ the local block) per KV group. Our port ran full causal attention (exact <=2048 ctx, divergent + O(ctx) beyond). This adds the real MSA. Converter: keep + quantize the 228 indexer tensors (index_{q,k}_proj int8, index_{q,k}_norm f32) instead of dropping them; --indexer mode gains M3 support for a supplemental add-on pass (no full re-convert). The sparse config already rides the flattened container config. Engine: reuses the DSA index-key cache (Ic) + index_hd/index_nh/idx_type slots (M3's sparse layers == MoE layers). load_cfg reads sparse_attention_config (index dim 128, 4 heads, block 128, top-16, local 1); guards the GLM DSA index read from clobbering M3's fields. attention_gqa: the indexer (Gemma-norm + same partial-ROT rope as the main attention, idx_k mirrored to Ic) runs per sparse layer, block-scores in f64, and picks top-k blocks (local forced in; greedy ties -> lowest index). The main core then attends only the selected blocks. VK attention is skipped on sparse layers. Validated on a tiny indexer-bearing checkpoint (block_size 4 so 24 tokens span 6 blocks): engine vs numpy oracle bit-exact — PREFILL 24/24 + incremental DECODE 20/20 (IDOT=0). The 2 near-tie mismatches during bring-up were a config bug (JSON-bool use_sparse_attention + the unguarded GLM index read zeroing index_hd), now fixed. --- c/colibri.c | 104 +++++++++++++++++++++++++++++++-- c/tools/convert_fp8_to_int4.py | 13 +++-- c/tools/make_m3tiny.py | 19 +++++- c/tools/oracle_m3.py | 42 ++++++++++++- 4 files changed, 162 insertions(+), 16 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 6e88b6167..dd6f73fd0 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -175,8 +175,9 @@ typedef struct { int first_dense, q_lora, kv_lora, qk_nope, qk_rope, qk_head, v_head, n_shared, vocab; int n_group, topk_group, norm_topk; int stop_ids[8], n_stop; /* eos_token_id dal config (GLM-5.2 ne ha 3!) */ - int index_topk, index_nh, index_hd; /* DSA lightning indexer */ - int8_t idx_type[128]; /* per layer: 1=full (calcola), 0=shared (riusa) */ + int index_topk, index_nh, index_hd; /* DSA lightning indexer (GLM); M3 MSA reuses index_nh/index_hd + Ic */ + int8_t idx_type[128]; /* per layer: 1=full (calcola), 0=shared (riusa); M3: sparse-attn layer */ + int msa, msa_blk, msa_topk_blk, msa_local_blk; /* M3 MSA: on, block size, top-k blocks, local blocks */ float eps, theta, attn_scale, routed_scale; int arch; /* ARCH_GLM | ARCH_M3 */ int n_kv_heads, head_dim, rotary; /* M3: GQA KV heads, head dim, partial-rope dims */ @@ -381,6 +382,8 @@ typedef struct { QT q_a, q_b, kv_a, kv_b, o; float *q_a_ln, *kv_a_ln; /* GQA (ARCH_M3): plain projections + per-head QK-norm weights [head_dim] */ QT q_p, k_p, v_p; float *q_hn, *k_hn; + /* MSA Lightning Indexer (ARCH_M3 sparse layers): block-selection branch */ + QT idx_q, idx_k; float *idx_qn, *idx_kn; #ifdef COLI_CUDA ColiCudaTensor *kv_b_shard[COLI_CUDA_MAX_DEVICES]; int shard_h0[COLI_CUDA_MAX_DEVICES],shard_hn[COLI_CUDA_MAX_DEVICES],n_kv_b_shard; @@ -1699,6 +1702,19 @@ static void load_cfg(Cfg *c, const char *snap){ { jval *fr=json_get(r,"moe_layer_freq"); c->first_dense=0; if(fr && fr->t==J_ARR){ int i=0; while(ilen && (int)fr->kids[i]->num==0) i++; c->first_dense=i; } } + /* MSA (Lightning Indexer): block-sparse attention on the non-dense layers. Reuses + * the DSA cache slots (index_hd/index_nh/idx_type/Ic). Sparse layers == MoE layers + * (sparse_attention_freq mirrors moe_layer_freq), so idx_type = (li >= first_dense). */ + { jval *sac=json_get(r,"sparse_attention_config"); + if(sac && gi(sac,"sparse_index_dim")>0){ /* use_sparse_attention is a JSON bool (gi reads num) */ + c->msa=1; + c->index_hd=gi(sac,"sparse_index_dim"); + c->index_nh=gi(sac,"sparse_num_index_heads"); + c->msa_blk=gi(sac,"sparse_block_size"); + c->msa_topk_blk=gi(sac,"sparse_topk_blocks"); + c->msa_local_blk=gi(sac,"sparse_local_block"); + for(int i=0;in_layers && i<128;i++) c->idx_type[i]=(i>=c->first_dense); + } } c->norm_topk=1; /* M3 always renormalizes the top-k weights */ c->n_group=1; /* no grouped routing (key absent -> 0) */ /* GQA rows ride the MLA cache aliases: Lc rows = K (n_kv*hd), Rc rows = V. */ @@ -1744,6 +1760,7 @@ static void load_cfg(Cfg *c, const char *snap){ fclose(gf); } } /* DSA lightning indexer: parametri + tipo per-layer (lista esplicita o formula freq/offset) */ + if(c->arch!=ARCH_M3){ /* GLM DSA indexer; M3 MSA set its own above */ c->index_topk=gi(r,"index_topk"); c->index_nh=gi(r,"index_n_heads"); c->index_hd=gi(r,"index_head_dim"); { jval *it=json_get(r,"indexer_types"); int freq=gi(r,"index_topk_freq"); if(freq<1) freq=1; @@ -1753,6 +1770,7 @@ static void load_cfg(Cfg *c, const char *snap){ c->idx_type[i] = !strcmp(it->kids[i]->str,"full"); else { int v=i-off+1; if(v<0) v=0; c->idx_type[i] = (v%freq)==0; } } } + } if(c->arch==ARCH_M3){ c->qk_head=c->head_dim; c->attn_scale = 1.f / sqrtf((float)c->head_dim); @@ -2405,6 +2423,12 @@ static void model_init_range(Model *m, const char *snap, int cap, l->o = qt_load(m,P("self_attn.o_proj.weight"), D, H*c->head_dim, dbits); l->q_hn = ld(m,P("self_attn.q_norm.weight")); l->k_hn = ld(m,P("self_attn.k_norm.weight")); + if(c->msa && i>=c->first_dense){ /* MSA Lightning Indexer (sparse layers) */ + l->idx_q = qt_load(m,P("self_attn.index_q_proj.weight"), c->index_nh*c->index_hd, D, dbits); + l->idx_k = qt_load(m,P("self_attn.index_k_proj.weight"), c->index_hd, D, dbits); + l->idx_qn = ld(m,P("self_attn.index_q_norm.weight")); + l->idx_kn = ld(m,P("self_attn.index_k_norm.weight")); + } } else { l->q_a = qt_load(m,P("self_attn.q_a_proj.weight"), c->q_lora, D, dbits); l->q_a_ln= ld(m,P("self_attn.q_a_layernorm.weight")); @@ -4216,6 +4240,15 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po matmul_qt(q,x,&l->q_p,S); matmul_qt(k,x,&l->k_p,S); matmul_qt(v,x,&l->v_p,S); + /* MSA Lightning Indexer (sparse layers): a small IDX_H-head scoring branch that picks + * the top-k key BLOCKS per query, so the main attention only reads those blocks. idx_q/ + * idx_k are roped like the main attention (same partial-ROT split-half NEOX); idx_k is + * mirrored into the persistent Ic cache. IDX_H == NK (one selection per KV group). */ + int IDX_H=c->index_nh, IDX_D=c->index_hd, BLK=c->msa_blk, TOPK=c->msa_topk_blk, LOCAL=c->msa_local_blk; + int is_sparse = (c->msa && layern_layers && c->idx_type[layer] && m->kv && m->kv->Ic); + float *iq=NULL, *ik=NULL; int *sel=NULL; + if(is_sparse){ iq=falloc((int64_t)S*IDX_H*IDX_D); ik=falloc((int64_t)S*IDX_D); + matmul_qt(iq,x,&l->idx_q,S); matmul_qt(ik,x,&l->idx_k,S); } m->t_aproj+=now_s()-tp0; for(int s=0;skv; @@ -4227,6 +4260,44 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po rms_head_g(kh,l->k_hn,hd,c->eps); rope_half_neox(kh,rot,pos,c->theta); } memcpy(ks->Lc[layer]+(int64_t)pos*KVd, ksr, (size_t)KVd*sizeof(float)); memcpy(ks->Rc[layer]+(int64_t)pos*KVd, vsr, (size_t)KVd*sizeof(float)); + if(is_sparse){ /* index heads: Gemma-norm + rope, mirror idx_k */ + float *iqs=iq+(int64_t)s*IDX_H*IDX_D, *iks=ik+(int64_t)s*IDX_D; + for(int h=0;hidx_qn,IDX_D,c->eps); rope_half_neox(ih,rot,pos,c->theta); } + rms_head_g(iks,l->idx_kn,IDX_D,c->eps); rope_half_neox(iks,rot,pos,c->theta); + memcpy(ks->Ic[layer]+(int64_t)pos*IDX_D, iks, (size_t)IDX_D*sizeof(float)); + } + } + if(is_sparse){ /* block selection: top-k blocks per (s, index head) */ + sel=malloc((size_t)S*IDX_H*TOPK*sizeof(int)); + for(int s=0;skv; + int pos=positions?positions[s]:pos_base+s, st0=ks->kv_start[layer]; + const float *Ic0=ks->Ic[layer]; int nblk=pos/BLK+1; + double *bscore=malloc((size_t)nblk*sizeof(double)); float *iqs=iq+(int64_t)s*IDX_H*IDX_D; + for(int hi=0;hibscore[b]) bscore[b]=d; + } + int qb=pos/BLK; /* local blocks always selected */ + for(int lb=0;lb=0) bscore[bb]=1e300; } + int topk = TOPK lowest block index */ + int best=-1; double bv=-1e300; + for(int b=0;bbv){ bv=bscore[b]; best=b; } } + sg[j]=best; + } + } + free(bscore); + } } double tc0=now_s(); int vk_core=0, vk_projected=0; (void)vk_core; (void)vk_projected; @@ -4237,7 +4308,7 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po * invalidated like the CUDA/MLA shadow on rewrite/rebind/resize). Single-sequence * decode only (no ragged mux); falls back to the CPU core on any failure. The * fused variant keeps ctx on-device and runs the o-projection in the same submit. */ - if(g_vk_attn && !kvs && !positions && S<=4 && layern_layers && + if(g_vk_attn && !is_sparse && !kvs && !positions && S<=4 && layern_layers && m->vk_kv_valid && m->kv->Lc[layer] && m->kv->Rc[layer]){ int st0=m->kv_start[layer], T=pos_base+S; if(T<=m->max_t && coli_vk_kv_ensure(layer,m->max_t,KVd,KVd)){ @@ -4271,6 +4342,28 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po const float *qh=qs+(int64_t)h*hd; int g=h/G; /* repeat_kv: query head h reads KV head h/G */ float *ah=att+(int64_t)h*win; + float *ch=cs+(int64_t)h*hd; + if(is_sparse){ /* MSA: attend only the selected key blocks of group g */ + int *sg=&sel[((int64_t)s*IDX_H+g)*TOPK]; + float mx=-1e30f; int nk=0; + for(int j=0;jT)t1=T; + for(int t=t0;tattn_scale; ah[nk++]=d; if(d>mx) mx=d; } } + float ssum=0; for(int t=0;tT)t1=T; + for(int t=t0;to,S); /* VK fused path already filled out */ m->t_aout+=now_s()-to0; - free(q); free(k); free(v); free(ctx); + free(q); free(k); free(v); free(ctx); free(iq); free(ik); free(sel); m->t_attn += now_s()-ta0; } @@ -7266,7 +7358,7 @@ static void kv_alloc(Model *m, int max_t){ free(k->Lc8); free(k->Rc8); free(k->Lsc); free(k->Rsc); k->Lc8=k->Rc8=NULL; k->Lsc=k->Rsc=NULL; } if(k->Ic){ for(int i=0;in_layers;i++) free(k->Ic[i]); free(k->Ic); k->Ic=NULL; } - if(m->has_dsa){ + if(m->has_dsa || c->msa){ /* index-key cache: GLM DSA or M3 MSA */ k->Ic=calloc(c->n_layers,sizeof(float*)); for(int i=0;in_layers;i++) if(c->idx_type[i]) k->Ic[i]=falloc((int64_t)max_t*c->index_hd); } diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 09c622c25..095dcbb45 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -202,8 +202,12 @@ def classify(name, n_layers, keep_mtp=False, keep_idx=False): if name.endswith((".weight_scale", ".weight_scale_2", ".input_scale")): return "consumed" li = layer_idx(name) if keep_idx: - # modalita' --indexer: SOLO i pesi del DSA lightning indexer dei layer principali - if li < 0 or li >= n_layers or "indexer" not in name: return "skip" + # modalita' --indexer: SOLO i pesi dell'indexer di scoring dei layer principali. + # GLM DSA usa "indexer"; MiniMax-M3 MSA usa "self_attn.index_" (Lightning Indexer). + # EN: --indexer mode: ONLY the scoring-indexer weights. GLM DSA = "indexer", + # MiniMax-M3 MSA = "self_attn.index_". + is_idx = ("indexer" in name) or (".index_" in name) + if li < 0 or li >= n_layers or not is_idx: return "skip" if name.endswith("norm.weight"): return "f32" return "q" # int8 consigliato (--ebits 8): pesi di scoring if keep_mtp: @@ -244,13 +248,12 @@ def classify(name, n_layers, keep_mtp=False, keep_idx=False): # .block_sparse_moe. -> .mlp. (router gate.weight + bias line up) # .experts.E.{w1,w3,w2}.weight -> .experts.E.{gate_proj,up_proj,down_proj}.weight # vision_tower / multi_modal_projector / patch_merge_mlp -> dropped (text-only port) -# .self_attn.index_* -> dropped (MSA index branch; separate pass later, -# like the GLM DSA indexer) +# .self_attn.index_{q,k}_{proj,norm} -> KEPT (MSA Lightning Indexer: q/k_proj int8, +# q/k_norm f32; used for block-sparse selection) # Returns the container name, or None to skip the tensor entirely. def m3_name(name): if name.startswith(("vision_tower.", "multi_modal_projector.", "patch_merge_mlp.")): return None - if ".self_attn.index_" in name: return None if name.startswith("language_model."): name = name[len("language_model."):] name = name.replace(".block_sparse_moe.e_score_correction_bias", ".mlp.gate.e_score_correction_bias") # the C loader's GLM name diff --git a/c/tools/make_m3tiny.py b/c/tools/make_m3tiny.py index dcaef72f2..5a8903f9d 100644 --- a/c/tools/make_m3tiny.py +++ b/c/tools/make_m3tiny.py @@ -2,7 +2,10 @@ engine-vs-oracle validation — the M3 counterpart of glm_tiny. Layer 0 dense + layer 1 MoE, GQA 4/2 heads with head_dim 8 (deliberately != hidden/heads to exercise the explicit head_dim path), partial rotary 4 of 8, all 4 experts present, nonzero router -bias (exercises the choice-vs-weight distinction). Convert with: +bias (exercises the choice-vs-weight distinction). The MoE/sparse layer also carries an +MSA Lightning Indexer (index_dim 8, 2 index heads = NK, block_size 4, top-2 blocks + +1 local) so the 24-token oracle sequence spans 6 blocks and exercises block selection. +Convert with: python3 tools/convert_fp8_to_int4.py --indir m3tiny --outdir m3tiny_i8 --ebits 8 --io-bits 8 then validate with tools/oracle_m3.py + the engine's REF/TF gate.""" import json, sys, os @@ -15,6 +18,7 @@ D, L, H, NK, HD = 16, 2, 4, 2, 8 ROT, E, TOPK, MI, DI, SI, V = 4, 4, 2, 8, 12, 8, 32 +IDX_DIM, IDX_H, BLK, TOPK_BLK, LOCAL_BLK = 8, NK, 4, 2, 1 # MSA Lightning Indexer t = {} def w(n, o, i): t[n] = (torch.randn(o, i) * 0.3).to(torch.bfloat16) @@ -34,6 +38,11 @@ def vv(n, o): t[n] = (torch.randn(o) * 0.3).to(torch.bfloat16) w(Lp+"self_attn.o_proj.weight", D, H*HD) vv(Lp+"self_attn.q_norm.weight", HD) vv(Lp+"self_attn.k_norm.weight", HD) + if i >= 1: # sparse layers carry the MSA Lightning Indexer + w(Lp+"self_attn.index_q_proj.weight", IDX_H*IDX_DIM, D) + w(Lp+"self_attn.index_k_proj.weight", IDX_DIM, D) # single index-k head + vv(Lp+"self_attn.index_q_norm.weight", IDX_DIM) + vv(Lp+"self_attn.index_k_norm.weight", IDX_DIM) if i == 0: # dense layer w(Lp+"mlp.gate_proj.weight", DI, D) w(Lp+"mlp.up_proj.weight", DI, D) @@ -64,7 +73,13 @@ def vv(n, o): t[n] = (torch.randn(o) * 0.3).to(torch.bfloat16) "rms_norm_eps": 1e-6, "rope_theta": 5000000, "vocab_size": V, "use_qk_norm": True, "qk_norm_type": "per_head", "use_gemma_norm": True, "hidden_act": "swigluoai", "swiglu_alpha": 1.702, "swiglu_limit": 7.0, - "eos_token_id": V-1}, + "eos_token_id": V-1, + "sparse_attention_config": { + "use_sparse_attention": True, "sparse_index_dim": IDX_DIM, + "sparse_num_index_heads": IDX_H, "sparse_block_size": BLK, + "sparse_topk_blocks": TOPK_BLK, "sparse_local_block": LOCAL_BLK, + "sparse_init_block": 0, "sparse_score_type": "max", + "sparse_attention_freq": [0, 1]}}, # layer 0 dense/full, layer 1 sparse "vision_config": {"hidden_size": 8}} json.dump(cfg, open(os.path.join(OUT, "config.json"), "w"), indent=1) print(f"{OUT}: {len(t)} tensors, D={D} L={L} H={H}/{NK} hd={HD} rot={ROT} E={E} top{TOPK} V={V}") diff --git a/c/tools/oracle_m3.py b/c/tools/oracle_m3.py index 930addceb..fb864b150 100644 --- a/c/tools/oracle_m3.py +++ b/c/tools/oracle_m3.py @@ -12,7 +12,7 @@ (gate=min(g,lim), up=clamp(±lim), (up+1)*gate*sigmoid(alpha*gate)); router sigmoid -> +bias for CHOICE only, raw-sigmoid weights renormalized, routed_out * routed_scaling + shared_out; pre-norm residual layers.""" -import json, sys, glob +import json, sys, glob, os import numpy as np from safetensors import safe_open @@ -29,6 +29,11 @@ EPS, TH = cfg["rms_norm_eps"], cfg["rope_theta"] A, LIM, RS = cfg["swiglu_alpha"], cfg["swiglu_limit"], cfg["routed_scaling_factor"] FD = sum(1 for x in (cfg.get("moe_layer_freq") or []) if x == 0) if cfg.get("moe_layer_freq") else 0 +SP = cfg.get("sparse_attention_config") or {} # MSA Lightning Indexer +SPARSE = bool(SP.get("use_sparse_attention")) +IDX_DIM, IDX_H = SP.get("sparse_index_dim"), SP.get("sparse_num_index_heads") +BLK, TOPK_BLK = SP.get("sparse_block_size"), SP.get("sparse_topk_blocks") +LOCAL_BLK = SP.get("sparse_local_block", 1) def deq(name): """Dequant a container tensor exactly as the engine reads it (int8 per-row, f32).""" @@ -78,13 +83,44 @@ def sigmoid(x): return 1.0 / (1.0 + np.exp(-x)) for h in range(H): q[s, h] = rope(rms(q[s, h], qn), s) for h in range(NK): k[s, h] = rope(rms(k[s, h], kn), s) G = H // NK + # --- MSA Lightning Indexer: per-query block selection on sparse layers (li >= FD) --- + # A small IDX_H-head dot-product branch scores every causal key, max-pools the scores + # into BLK-sized blocks, forces the LOCAL_BLK current/preceding blocks in, and keeps the + # top-TOPK_BLK blocks per index head. Index head hi == KV group g (IDX_H == NK). Uses the + # SAME partial-ROT split-half rope as the main attention (index_head_dim > ROT). + sel = None + if SPARSE and li >= FD: + iq = nrm @ deq(P + "self_attn.index_q_proj.weight").T # [S, IDX_H*IDX_DIM] + ik = nrm @ deq(P + "self_attn.index_k_proj.weight").T # [S, IDX_DIM] (single k head) + iqn = T[P + "self_attn.index_q_norm.weight"].astype(np.float32) + ikn = T[P + "self_attn.index_k_norm.weight"].astype(np.float32) + iq = iq.reshape(S, IDX_H, IDX_DIM) + for s in range(S): + for h in range(IDX_H): iq[s, h] = rope(rms(iq[s, h], iqn), s) + ik[s] = rope(rms(ik[s], ikn), s) + sel = [[None] * IDX_H for _ in range(S)] + for s in range(S): + for hi in range(IDX_H): + scores = iq[s, hi].astype(np.float64) @ ik[:s+1].astype(np.float64).T # [s+1] causal + nblk = s // BLK + 1 + bscore = np.full(nblk, -np.inf) + for t in range(s + 1): + b = t // BLK + if scores[t] > bscore[b]: bscore[b] = scores[t] + qb = s // BLK # local boost -> always selected + for l in range(LOCAL_BLK): + if qb - l >= 0: bscore[qb - l] = np.inf + topk = min(TOPK_BLK, nblk) + order = np.argsort(-bscore, kind="stable") # ties -> lower block index first + sel[s][hi] = set(int(b) for b in order[:topk]) ctx = np.zeros((S, H, HD), np.float32) for s in range(S): for h in range(H): g = h // G - sc = (q[s, h] @ k[:s+1, g].T) / np.sqrt(HD) + kk = np.array([t for t in range(s+1) if sel is None or (t // BLK) in sel[s][g]]) + sc = (q[s, h] @ k[kk, g].T) / np.sqrt(HD) sc = sc - sc.max(); e = np.exp(sc); a = e / e.sum() - ctx[s, h] = a @ v[:s+1, g] + ctx[s, h] = a @ v[kk, g] x = x + ctx.reshape(S, H * HD) @ deq(P + "self_attn.o_proj.weight").T nrm = rms(x, T[P + "post_attention_layernorm.weight"].astype(np.float32)) if li < FD: # dense From 453507e36b3cbe114213caceb698ef1f0a801e20 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 23 Jul 2026 16:30:30 +0200 Subject: [PATCH 08/22] =?UTF-8?q?engine:=20MSA=20auto-detect=20=E2=80=94?= =?UTF-8?q?=20fall=20back=20to=20full=20attention=20when=20indexer=20weigh?= =?UTF-8?q?ts=20absent?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A container converted before the indexer was kept still announces sparse_attention_ config, so the loader would try (and fail) to load index_q_proj on every sparse layer. Probe the first sparse layer's indexer weight; if absent, disable MSA (full causal attention, exact <=2048 ctx) and point the user at the --indexer pass. Same auto-detect discipline as the DSA indexer / MTP. --- c/colibri.c | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/c/colibri.c b/c/colibri.c index dd6f73fd0..b3233c690 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2409,6 +2409,14 @@ static void model_init_range(Model *m, const char *snap, int cap, m->eheat=calloc(NR,sizeof(uint32_t*)); m->elast=calloc(NR,sizeof(uint32_t*)); m->elast_dc=calloc(NR,sizeof(uint32_t*)); m->elast_pre=calloc(NR,sizeof(uint32_t*)); + /* MSA auto-detect: the config announces sparse attention, but a container converted + * before the indexer was kept has no index weights — fall back to full attention + * (exact <=2048 ctx) rather than failing the load. Add them with a `--indexer` pass. */ + if(c->msa){ char inm[300]; + snprintf(inm,sizeof(inm),"model.layers.%d.self_attn.index_q_proj.weight",c->first_dense); + if(!st_has(&m->S,inm)){ c->msa=0; for(int i=0;in_layers;i++) c->idx_type[i]=0; + fprintf(stderr,"[MSA] indexer weights absent — full causal attention " + "(add them with: convert --arch m3 --indexer for block-sparse long context)\n"); } } m->kv=calloc(1,sizeof(KVState)); m->kv_start=m->kv->kv_start=calloc(NR,sizeof(int)); for(int i=layer_begin;i Date: Thu, 23 Jul 2026 16:48:57 +0200 Subject: [PATCH 09/22] convert: enable the --indexer supplemental pass for MiniMax-M3 MSA MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three fixes so a container converted before the indexer was kept can get it added without a full 34-min re-convert: - drop the '--arch m3 --indexer not implemented' guard (only --mtp is refused now); - the shard pre-filter matched GLM's 'indexer' name only — also match M3's 'self_attn.index_' so the right shards are read; - layer_idx() only parsed 'model.layers.N...'; find 'layers.N' anywhere so it also handles M3's 'language_model.model.layers.N...' VL prefix. Extracts the 228 indexer weights (+scales) to out-idx-*.safetensors; drop them beside the main container (st.h scans any *.safetensors). Applied to the box container + mirror: MSA now active on the real model (no [MSA] fallback, output correct). --- c/tools/convert_fp8_to_int4.py | 22 ++++++++++++++-------- 1 file changed, 14 insertions(+), 8 deletions(-) diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 095dcbb45..cf1112ff1 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -189,10 +189,13 @@ def quant_int2(w, bits): # -> (qbytes U8 [O*ceil(I/4)], s # ---------- classificazione dei tensori ---------- def layer_idx(name): + # find "layers." anywhere: handles both "model.layers.N..." (mapped) and the raw + # "language_model.model.layers.N..." (MiniMax-M3 VL prefix) forms. p = name.split(".") - if len(p) > 2 and p[0] == "model" and p[1] == "layers": - try: return int(p[2]) - except ValueError: return -1 + for i in range(len(p) - 1): + if p[i] == "layers": + try: return int(p[i + 1]) + except ValueError: return -1 return -1 def classify(name, n_layers, keep_mtp=False, keep_idx=False): @@ -622,9 +625,11 @@ def main(): print(f"[ARCH] auto-detected: {a.arch}" + (f" (model_type {mt})" if mt else "")) if a.n_layers is None: a.n_layers = 60 if a.arch == "m3" else 78 - if a.arch == "m3" and (a.mtp or a.indexer): - raise SystemExit("--arch m3: the checkpoint ships no MTP tensors, and the MSA index " - "branch pass is not implemented yet (dropped in the main pass, like DSA)") + if a.arch == "m3" and a.mtp: + raise SystemExit("--arch m3: the checkpoint ships no MTP tensors") + # --arch m3 --indexer: supplemental pass for the MSA Lightning Indexer weights + # (self_attn.index_{q,k}_{proj,norm}) -> out-idx-*.safetensors, kept next to a main + # container that predates the indexer being kept in the main pass. # 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' @@ -730,8 +735,9 @@ def get_slice(s, n): return None 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} + else: # GLM DSA "indexer" or MiniMax-M3 MSA "self_attn.index_" + want = {v for k, v in wmap.items() + if ("indexer" in k or ".index_" 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 From ae605a35b86fe1e0d02a6cc121efdf3bf4db5b5f Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Fri, 24 Jul 2026 11:23:55 +0200 Subject: [PATCH 10/22] kv-persist: save/restore the MSA index-key cache (Ic) for M3 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit kv_persist.h gated the Ic rows on has_dsa (GLM DSA), which is false for M3 (index_topk=0), so .coli_kv persisted K/V but not the MSA index keys: a resumed conversation carried uninitialized Ic for every restored position, silently corrupting block selection once the conversation passed 2048 tokens (topk_blocks * block_size). Gate on (has_dsa || msa) — the same condition kv_alloc uses for Ic — in the header, record size, append and load. Old M3 .coli_kv files now mismatch the header (h[3] 0 -> index_hd) and take the existing 'different model or version' reject path (start over, no misparse). Validated on the real 225 GB container: - record size 245764 -> 274948 B/token (= +57 sparse layers x 128 f32) - old-format file: clean reject + start-over - end-to-end: a 2363-token needle document (needle in block 15 of 19) saved by one process; a FRESH process resumed 2365 tokens (0.1 s, no re-prefill) and answered the needle question exactly. Without the fix the zero-key tie-break selects blocks {0..14, local} and drops block 15, so this recall only works through the restored index keys. --- c/colibri.c | 2 +- c/kv_persist.h | 13 +++++++------ 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index b3233c690..4df8b09f4 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -8699,7 +8699,7 @@ static void repin_pass_limit(Model *m,int limit){ /* ---- KV SU DISCO: la conversazione si riapre CALDA (KVSAVE=0 disattiva) ---- * Il re-prefill di una chat riaperta costa ore su questo disco; la KV compressa MLA * costa ~182 KB/token. File /.coli_kv append-only: header (magic + dimensioni + - * nrec) e un record per posizione [tok i32][Lc+Rc dei 78 layer][Ic DSA]. A fine turno + * nrec) e un record per posizione [tok i32][Lc+Rc dei 78 layer][Ic DSA/MSA]. A fine turno * si appendono SOLO le posizioni nuove e si riscrive nrec per ultimo: un crash a meta' * append lascia nrec vecchio = file coerente. La riga KV del layer MTP non si salva: * al resume kv_start=-1 e la finestra di draft riparte da sola. */ diff --git a/c/kv_persist.h b/c/kv_persist.h index 2ad5ffb7b..70fbd8dd6 100644 --- a/c/kv_persist.h +++ b/c/kv_persist.h @@ -1,6 +1,7 @@ /* kv_persist.h — .coli_kv on-disk KV cache persistence. - * Conversations reopen warm across engine restarts: the compressed MLA KV-cache - * is appended incrementally after every turn, crash-safe (nrec written last). + * Conversations reopen warm across engine restarts: the KV cache (MLA rows, or + * GQA K/V riding the same aliases) plus the DSA/MSA index-key cache Ic is + * appended incrementally after every turn, crash-safe (nrec written last). * Include after Model/KVState/Cfg are defined; requires now_s() and g_draft. */ #ifndef KV_PERSIST_H #define KV_PERSIST_H @@ -15,7 +16,7 @@ static void kv_hdr(Model *m, int32_t *h, int nrec){ Cfg *c=&m->c; int nic=0; for(int i=0;in_layers;i++) if(m->Ic && m->Ic[i]) nic++; h[0]=c->n_layers; h[1]=c->kv_lora; h[2]=c->qk_rope; - h[3]=m->has_dsa?c->index_hd:0; h[4]=nic; h[5]=c->vocab; h[6]=nrec; + h[3]=(m->has_dsa||c->msa)?c->index_hd:0; h[4]=nic; h[5]=c->vocab; h[6]=nrec; h[7]=g_tq?((g_tq_codec<<8)|g_tq_bits):(g_kv8?1:0); /* format tag: 0=f32, 1=kv8; TQ: codec<<8 | bit width */ } @@ -24,7 +25,7 @@ static int64_t kv_rec_bytes(Model *m){ int64_t rec = 4 + (g_tq ? (int64_t)c->n_layers*(coli_kvq_row_bytes(c->kv_lora,g_tq_bits,g_tq_codec)+coli_kvq_row_bytes(c->qk_rope,g_tq_bits,g_tq_codec)+8) : g_kv8 ? (int64_t)c->n_layers*(c->kv_lora+c->qk_rope+8) : (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4); - if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4; + if(m->has_dsa||c->msa) for(int i=0;in_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4; return rec; } @@ -100,7 +101,7 @@ static void kv_disk_append(Model *m, const int *hist, int len){ memcpy(b, m->Rc[i]+(int64_t)p*c->qk_rope,(size_t)c->qk_rope*4); b+=c->qk_rope*4; } } - if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]){ + if(m->has_dsa||c->msa) for(int i=0;in_layers;i++) if(m->Ic[i]){ memcpy(b, m->Ic[i]+(int64_t)p*c->index_hd, (size_t)c->index_hd*4); b+=c->index_hd*4; } fwrite(k->disk_buf, 1, (size_t)rec, f); @@ -197,7 +198,7 @@ static int kv_disk_load(Model *m, int *hist, int maxctx){ fread(m->Rc[i]+(int64_t)p*c->qk_rope, 4, c->qk_rope, f)!=(size_t)c->qk_rope){ nrec=p; goto out; } } } - if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) + if(m->has_dsa||c->msa) for(int i=0;in_layers;i++) if(m->Ic[i]) if(fread(m->Ic[i]+(int64_t)p*c->index_hd, 4, c->index_hd, f)!=(size_t)c->index_hd){ nrec=p; goto out; } } out: From b295ffee426622fb722b4eb7c05342eac6a69c8a Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Fri, 24 Jul 2026 11:58:18 +0200 Subject: [PATCH 11/22] vulkan: generic VK dense matmul for resident weights (COLI_VK_DENSE, M3) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit M3's GQA has no MLA low-rank bottleneck: q_proj and o_proj are 50 MB of int8 each across 60 layers, ~6 GB of weight reads EVERY token that rode the same ~28 GB/s DDR4 the routed experts stream over — the reason a 428B/23B-active model barely beat the 744B GLM (whose attention/dense lives in VRAM via CUDA_DENSE). matmul_qt_ex had Metal and CUDA branches but no Vulkan one; COLI_VK_DENSE only wired GLM's MLA call sites. Add a generic VK branch to matmul_qt_ex behind COLI_VK_DENSE=1, gated on a per-tensor vk_gemm flag set at load for M3's resident weights (q/k/v/o, indexer, dense MLP; shared experts already had wired sites) — routed expert QTs are slab-transient and never marked. At decode only tensors >= COLI_VK_GEMM_MB (default 8) are worth the per-call submit; a batched forward (S>=8) offloads every marked tensor. Weights upload lazily at the dense priority class (0.75 > tier 0.4), so the launcher reserve grows to 8 GB and the expert tier self-sizes into the rest. Measured (RX 9070, 512->~170-200-expert tier, 2-drive mirror, ngen 64): decode topp 0.7: 2.12-2.20 -> 2.6-2.84 tok/s (+25-30%) decode topp 0.5: 2.49 -> 3.0-3.10 tok/s prefill 516 tok: 152.7 -> 102.3 s (-33%); attention projections -67% Output verified (greedy Paris + coherent lighthouse prose). Launchers updated to COLI_VK_DENSE=1 RESERVE=8 and the stale MSA caveat rewritten. --- c/colibri.c | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/c/colibri.c b/c/colibri.c index 4df8b09f4..01570f85f 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -275,6 +275,8 @@ typedef struct { * fault fails at S=1, records 1, and is never retried -- the old behaviour, reached * as a special case of the general rule rather than as a separate one. */ int cuda_fail_s; + int vk_gemm; /* resident tensor eligible for the generic VK dense matmul in + * matmul_qt_ex (marked at load; never set on reused expert slots) */ } QT; static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ int64_t n=(int64_t)t->O*t->I; @@ -618,6 +620,10 @@ static inline int vk_reg_served(int layer,int eid){ } static int g_vk_dense; /* COLI_VK_DENSE=1: run the resident dense matmuls (attention * projections + shared expert) on Vulkan too */ +static int g_vk_gemm_mb; /* COLI_VK_GEMM_MB (default 8): minimum resident-weight MB for + * the generic VK dense matmul at decode — the per-call submit + * only beats the CPU's DDR4 read on big tensors. A batched + * forward (S>=8) offloads every vk_gemm-marked tensor. */ static int g_vk_budget2; /* COLI_VK_EXPERTS2: dev2 expert-tier cap (with COLI_VK_DEV2) */ static int g_vk_reg_n2; /* experts resident on the dev2 tier */ /* Resolve the main shader path (#523): COLI_VK_SHADERS may be the qmatmul.spv file itself @@ -1173,6 +1179,16 @@ static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot) "narrower calls will still try the GPU\n", w->O,w->I,w->cuda_device,S); cuda_disabled_note(); } +#endif +#ifdef COLI_VULKAN + /* COLI_VK_DENSE=1: resident dense/attention weights (vk_gemm, marked at load) on the + * GPU — M3's GQA q/o projections alone read ~6 GB/token of int8 from DDR4 on the CPU + * path. Decode offloads only tensors >= COLI_VK_GEMM_MB (the per-call submit only pays + * for itself on big weight reads); a batched prefill (S>=8) offloads every marked + * tensor. Routed-expert QTs are slab-transient and never marked. */ + if(g_vulkan && g_vk_dense && w->vk_gemm && !spec_pinned() && !omp_in_parallel() + && (S>=8 || qt_bytes(w) >= (int64_t)g_vk_gemm_mb<<20) + && vk_matmul_qt(w,y,x,S)) return; #endif if(w->fmt==0){ matmul(y,x,w->qf,S,w->I,w->O); return; } if(w->fmt==4){ @@ -2436,7 +2452,9 @@ static void model_init_range(Model *m, const char *snap, int cap, l->idx_k = qt_load(m,P("self_attn.index_k_proj.weight"), c->index_hd, D, dbits); l->idx_qn = ld(m,P("self_attn.index_q_norm.weight")); l->idx_kn = ld(m,P("self_attn.index_k_norm.weight")); + l->idx_q.vk_gemm=l->idx_k.vk_gemm=1; } + l->q_p.vk_gemm=l->k_p.vk_gemm=l->v_p.vk_gemm=l->o.vk_gemm=1; /* COLI_VK_DENSE offload */ } else { l->q_a = qt_load(m,P("self_attn.q_a_proj.weight"), c->q_lora, D, dbits); l->q_a_ln= ld(m,P("self_attn.q_a_layernorm.weight")); @@ -2460,6 +2478,11 @@ static void model_init_range(Model *m, const char *snap, int cap, l->up_proj = qt_load(m,P("mlp.up_proj.weight"), c->dense_inter, D, dbits); l->down_proj = qt_load(m,P("mlp.down_proj.weight"), D, c->dense_inter, dbits); qt_planarize(&l->gate_proj); qt_planarize(&l->up_proj); qt_planarize(&l->down_proj); /* K1 */ + /* K1 and the VK dense path never claim the same bytes: planar_on() is + * hard-off on GPU builds (the backends read q4 in pair layout), so at + * most one of the two marks is live in any given binary. */ + if(c->arch==ARCH_M3) /* GLM keeps its wired MLA VK sites */ + l->gate_proj.vk_gemm=l->up_proj.vk_gemm=l->down_proj.vk_gemm=1; } else { l->router=ld(m,P("mlp.gate.weight")); l->router_bias=ld(m,P("mlp.gate.e_score_correction_bias")); @@ -11193,6 +11216,7 @@ int main(int argc, char **argv){ * ~6 GB tier + ~8 GB dense leaves headroom for the long-context KV mirror). */ g_vk_budget = getenv("COLI_VK_EXPERTS") ? atoi(getenv("COLI_VK_EXPERTS")) : 320; g_vk_dense = getenv("COLI_VK_DENSE") ? atoi(getenv("COLI_VK_DENSE")) : 0; + g_vk_gemm_mb = getenv("COLI_VK_GEMM_MB") ? atoi(getenv("COLI_VK_GEMM_MB")) : 8; g_vk_attn = getenv("COLI_VK_ATTN") ? atoi(getenv("COLI_VK_ATTN")) : 0; fprintf(stderr,"[VK] expert tier active: routed quantized experts on the GPU (budget %d)%s%s\n", g_vk_budget, g_vk_dense ? " + dense projections + shared expert" : "", From e1035debf592420b0d11424aa6778d426f7915c9 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Fri, 24 Jul 2026 12:15:05 +0200 Subject: [PATCH 12/22] vulkan/msa: remaining int4 decode+prefill levers for M3 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Four levers on top of the VK dense matmul (95a565d): - lm_head vk_gemm: the 1.2 GB int8 logit matmul rides COLI_VK_DENSE. - q+k pair: k_p's 3 MB joins q_p's submit via vk_matmul_pair_qt (a lone k submit costs more than its CPU read; paired it is free). - fused shared expert un-gated for M3: the gate_up shader takes the activation as a push constant since 8dd7fae, and fmt=4/g64 rides the same grouped path the expert tier uses — the !g_act_swigluoai and fmt gate predated both. 3 submits/layer -> 1 across 57 layers. Also set the activation BEFORE vk_registry_fill's early returns, so a tier-less run (COLI_VK_EXPERTS=0) can not dispatch silu on an swigluoai model. - MSA selection: t-outer loop (each cached key row read once for all 4 index heads), AVX2 4-lane f64 FMA scoring dot (the 64K-context decode hot spot), OMP over prefill rows (was single-threaded on 12 cores). Oracle gate after the selection rewrite: PREFILL 24/24 (exercises the OMP path), incremental DECODE 20/20 (IDOT=0, tiny checkpoint). Measured (RX 9070, R8 tier 168 + dense, 2-drive mirror, ngen 64): decode topp 0.5: 3.02-3.10 -> 3.49-3.52 tok/s decode topp 0.7: 2.78-2.84 (unchanged — 154 experts x 28.3 MB = 4.4 GB/token sits ON the dual-SSD read floor; compute offloads only pay below it, which is the standing argument for topp<=0.55 or int3) prefill 516 tok: 102.3 -> 96.8 s (152.7 before VK dense) RAM_GB=58 probe: no gain (hit 84.8->85.9, wash) — RAM budget is spent. Quality: Paris + coherent lighthouse prose at 0.5 and 0.7. --- c/colibri.c | 83 +++++++++++++++++++++++++++++++++++++++++------------ 1 file changed, 65 insertions(+), 18 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 01570f85f..ef7baab31 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2390,6 +2390,8 @@ static void model_init_range(Model *m, const char *snap, int cap, if(load_boundaries){ m->embed = qt_load(m,"model.embed_tokens.weight", c->vocab, D, io_bits); m->lm_head = qt_load(m,"lm_head.weight", c->vocab, D, io_bits); + if(c->arch==ARCH_M3) m->lm_head.vk_gemm=1; /* 1.2 GB int8 read EVERY token: prime + * COLI_VK_DENSE material (~28 -> ~3 ms) */ m->final_norm = ld(m,"model.norm.weight"); } m->L=calloc(c->n_layers,sizeof(Layer)); @@ -4260,6 +4262,31 @@ static void rms_head_g(float *v, const float *w, int n, float eps){ /* per-head float r=1.f/sqrtf((float)(ms/n)+eps); for(int i=0;ic; int H=c->n_heads, NK=c->n_kv_heads, hd=c->head_dim; @@ -4268,8 +4295,15 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po float *q=falloc((int64_t)S*H*hd), *k=falloc((int64_t)S*KVd), *v=falloc((int64_t)S*KVd); float *ctx=falloc((int64_t)S*H*hd); double tp0=now_s(); - matmul_qt(q,x,&l->q_p,S); - matmul_qt(k,x,&l->k_p,S); + int qk_vk=0; (void)qk_vk; +#ifdef COLI_VULKAN + /* q+k share x and one submit: k_p's 3 MB rides q_p's dispatch for free (a + * lone k submit would cost more than its CPU read, hence no solo offload). */ + if(g_vulkan && g_vk_dense && l->q_p.vk_gemm && !spec_pinned() && !omp_in_parallel() + && (S>=8 || qt_bytes(&l->q_p) >= (int64_t)g_vk_gemm_mb<<20)) + qk_vk = vk_matmul_pair_qt(&l->q_p,q,&l->k_p,k,x,S); +#endif + if(!qk_vk){ matmul_qt(q,x,&l->q_p,S); matmul_qt(k,x,&l->k_p,S); } matmul_qt(v,x,&l->v_p,S); /* MSA Lightning Indexer (sparse layers): a small IDX_H-head scoring branch that picks * the top-k key BLOCKS per query, so the main attention only reads those blocks. idx_q/ @@ -4301,21 +4335,27 @@ static void attention_gqa(Model *m, Layer *l, int layer, float *x, int S, int po } if(is_sparse){ /* block selection: top-k blocks per (s, index head) */ sel=malloc((size_t)S*IDX_H*TOPK*sizeof(int)); + #pragma omp parallel for schedule(static) if(S>4) /* rows independent; Ic rows <= pos + * were all mirrored above */ for(int s=0;skv; int pos=positions?positions[s]:pos_base+s, st0=ks->kv_start[layer]; const float *Ic0=ks->Ic[layer]; int nblk=pos/BLK+1; - double *bscore=malloc((size_t)nblk*sizeof(double)); float *iqs=iq+(int64_t)s*IDX_H*IDX_D; - for(int hi=0;hibscore[b]) bscore[b]=d; + double *bscore=malloc((size_t)IDX_H*nblk*sizeof(double)); float *iqs=iq+(int64_t)s*IDX_H*IDX_D; + for(int j=0;jbscore[(int64_t)hi*nblk+b]) bscore[(int64_t)hi*nblk+b]=d; } + } + for(int hi=0;hi=0) bscore[bb]=1e300; } + for(int lb=0;lb=0) bs[bb]=1e300; } int topk = TOPKbv){ bv=bscore[b]; best=b; } } + if(!taken && bs[b]>bv){ bv=bs[b]; best=b; } } sg[j]=best; } } @@ -6525,12 +6565,17 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int #endif if(!shared_cuda){ #ifdef COLI_VULKAN - /* Whole shared expert as ONE fused submit (gate+up+silu -> down, hidden on-device) + /* Whole shared expert as ONE fused submit (gate+up+act -> down, hidden on-device) * via the expert-group primitive with count=1 — replaces 3 separate VK matmuls - * (x read once, 1 fence instead of 3). Falls through to the per-matmul chain. */ + * (x read once, 1 fence instead of 3). The gate_up shader takes the activation + * as a push constant (silu/swigluoai, coli_vk_set_activation) so M3 qualifies; + * fmt=4 rides the same grouped-int4 path the expert tier uses. Falls through + * to the per-matmul chain on any mismatch. */ int fsh=l->sh_gate.fmt; - if(g_vk_dense && !g_act_swigluoai && !omp_in_parallel() && (fsh==1||fsh==2||fsh==5) && - l->sh_up.fmt==fsh && l->sh_down.fmt==fsh){ /* fused shader hardcodes silu — M3 uses the per-matmul chain */ + if(g_vk_dense && !omp_in_parallel() && + (fsh==1||fsh==2||fsh==5||(fsh==4 && l->sh_gate.gs>=8 && l->sh_gate.gs%8==0)) && + l->sh_up.fmt==fsh && l->sh_up.gs==l->sh_gate.gs && l->sh_down.fmt==fsh && + (fsh!=4 || (l->sh_down.gs>=8 && l->sh_down.gs%8==0))){ #define SW_(t) ((t).fmt==1?(const void*)(t).q8:(const void*)(t).q4) if(coli_vk_tensor_ensure(&l->sh_gate.vk,SW_(l->sh_gate),l->sh_gate.s,fsh,D,sI,l->sh_gate.gs)&& coli_vk_tensor_ensure(&l->sh_up.vk, SW_(l->sh_up), l->sh_up.s, fsh,D,sI,l->sh_up.gs)&& @@ -9584,11 +9629,13 @@ static void vk_dense_preload(Model *m){ static void vk_registry_fill(Model *m){ Cfg *c=&m->c; int E=c->n_experts, NL=c->n_layers; - if(!g_vulkan || g_vk_budget<=0) return; + if(!g_vulkan) return; /* The fused gate+up shader selects its activation via a push constant — silu for GLM, * swigluoai for MiniMax-M3 — so the expert tier runs on any of our arches. Set the - * model's choice once here, before the tier fills (all gate_up dispatches read it). */ + * model's choice BEFORE any early return: the fused shared-expert path dispatches + * gate_up too, even when the tier is empty or disabled (COLI_VK_EXPERTS=0). */ coli_vk_set_activation(g_act_swigluoai, c->swiglu_alpha, c->swiglu_limit); + if(g_vk_budget<=0) return; int64_t nz=0; for(int i=0;ieusage[i]) for(int e=0;eeusage[i][e]) nz++; if(!nz){ fprintf(stderr,"[VK] expert tier: no usage history yet — tier empty this run " From 844c8f45ebcceb403f342e94caf5c723e4d61212 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Fri, 24 Jul 2026 15:39:35 +0200 Subject: [PATCH 13/22] =?UTF-8?q?convert:=20--idx-bits=20=E2=80=94=20keep?= =?UTF-8?q?=20the=20MSA/DSA=20scoring-indexer=20at=20int8=20in=20a=20main?= =?UTF-8?q?=20pass?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Since the indexer weights are kept in the main pass (MSA), they fell into the generic resident class and silently followed --ebits. Block selection is discrete: quantization noise flips top-k choices rather than blurring them, and the validated real-model MSA config used int8 (supplemental pass). New 'idx' class + --idx-bits (default 8) pins them regardless of ebits/xbits; the tensors are tiny (~215 MB) so the cost is nil. Flag-path smoke on the tiny checkpoint: --ebits 4 --xbits 3 --group-size 64 converts (routed experts fmt=5 int3-g64, [MIXED] idx=8bit) and the engine loads and runs the result. --- c/tools/convert_fp8_to_int4.py | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index cf1112ff1..4e4f1a099 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -232,6 +232,11 @@ def classify(name, n_layers, keep_mtp=False, keep_idx=False): # "attn" = other attention projections (q_a, q_b, kv_a) # "dmlp" = dense MLP (first 3 layers) if "shared_experts" in name: return "sh" + # "idx" = MSA/DSA scoring-indexer projections (MiniMax-M3 self_attn.index_*). + # Selection is DISCRETE (top-k blocks): quantization noise flips choices rather + # than blurring them, so these tiny tensors (~215 MB) default to int8 instead + # of following ebits (the validated real-model MSA config used int8). + if ".index_" in name and name.endswith("proj.weight"): return "idx" if name.endswith("o_proj.weight"): return "o" if name.endswith("kv_b_proj.weight"): return "kvb" if any(name.endswith(k) for k in ("q_a_proj.weight", "q_b_proj.weight", @@ -512,6 +517,9 @@ def main(): help="bits for kv_b_proj (reconstructs KV cache on every decode). Default=ebits") ap.add_argument("--attn-bits", type=int, default=None, help="bits for other attention projections (q_a, q_b, kv_a). Default=ebits") + ap.add_argument("--idx-bits", type=int, default=8, + help="bits for the MSA/DSA scoring-indexer projections (block selection is " + "discrete — noise flips choices; tiny tensors, keep int8). Default=8") ap.add_argument("--dmlp-bits", type=int, default=None, help="bits for dense MLP (first 3 layers). Default=ebits") ap.add_argument("--group-size", type=int, default=64, @@ -599,6 +607,7 @@ def main(): if a.kvb_bits is not None: bits_map["kvb"] = a.kvb_bits if a.attn_bits is not None: bits_map["attn"] = a.attn_bits if a.dmlp_bits is not None: bits_map["dmlp"] = a.dmlp_bits + bits_map["idx"] = a.idx_bits # always set: default 8 (see --idx-bits) if bits_map: print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items()))) From 533be9ae82580832515fc5db5b803547227ea08b Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 13 Aug 2026 23:19:46 +0200 Subject: [PATCH 14/22] oracle_m3: gate invocation matches dev (SNAP=), document the IDOT=0 exact compare The docstring predates the rebase: dev's engine takes the model dir as SNAP= (COLI_MODEL is not an engine env upstream). Also spell out that the T/T expectation holds under IDOT=0 -- the exact f32 kernels reproduce this oracle to ~1e-6 max logit delta, while the default int8-activation IDOT path can flip the argmax at a borderline position of the random tiny fixture (measured: one 0.058-margin position at 23/24), which is activation quantization, not a port defect. --- c/tools/oracle_m3.py | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/c/tools/oracle_m3.py b/c/tools/oracle_m3.py index fb864b150..42d24f0c2 100644 --- a/c/tools/oracle_m3.py +++ b/c/tools/oracle_m3.py @@ -3,7 +3,10 @@ differences are implementation bugs, not quantization), teacher-forces a fixed token sequence with the pinned M3 conventions, and writes the engine's REF gate: ref_m3.json {prompt_ids, full_ids, tf_pred} + oracle_logits.npy [T,V] -Validate: REF=ref_m3.json TF=1 COLI_MODEL= ./colibri 8 (expect T/T) +Validate: SNAP= REF=ref_m3.json TF=1 TF_DECODE=1 ./colibri 8 +(expect T/T on both gates; run with IDOT=0 — the exact f32 kernels match this +oracle to ~1e-6, while the int8-activation IDOT default can flip an argmax at +a borderline position, which is quantization, not a port bug) Conventions (transformers modeling_minimax_m3_vl, pinned 2026-07-22): Gemma RMSNorm x/rms*(1+w) in f32; per-head QK-norm BEFORE partial NEOX RoPE From 2b62de21a0f43a4b577fa01e698f0dbb8a074ba9 Mon Sep 17 00:00:00 2001 From: Gerald Corzo Date: Sat, 15 Aug 2026 00:50:29 +0200 Subject: [PATCH 15/22] MiniMax-M3: COLI_MSA=0 kill-switch, M3/MSA banner+PROF, and tool-call rendering MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Three of the WIP gaps called out in #601: 1. COLI_MSA=0 A/B kill-switch (colibri.c). Force full causal attention even when the MSA indexer weights are present, by zeroing c->msa + idx_type the same two lines the indexer-absent auto-detect fallback uses — so it rides the already-tested full-attention path for measuring MSA's long-context win. 2. Banner/PROF cosmetics (colibri.c). The startup banner and the PROF attention-bound verdict now name MSA/M3 instead of hardcoding GLM/DSA (banner peeks config.json model_type before model_init; PROF branches on m->c.arch). 3. Tool calls for the MiniMax-M3 template (openai_server.py). Render the ]<]minimax[>[ XML dialect — the developer declaration, assistant tool_calls, and ]~b]tool … tool results — and parse it back, so render_chat no longer returns 400 for `tools` and parse_tool_calls dispatches to an M3 parser. Mirrors chat_template.jinja (to_xml recursion, nested objects/arrays, schema type coercion via _tool_param_types). Verified: MSA oracle still 24/24 + 20/20; COLI_MSA=0 flips to full attention; tool-call render/parse round-trips (nested objects/arrays + type coercion); full `make test` suite 167/167 OK (13 skipped). --- c/colibri.c | 28 ++++++- c/family_registry.py | 2 +- c/openai_server.py | 181 +++++++++++++++++++++++++++++++++++++++++-- 3 files changed, 201 insertions(+), 10 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index ef7baab31..6d1d7d1f2 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2435,6 +2435,13 @@ static void model_init_range(Model *m, const char *snap, int cap, if(!st_has(&m->S,inm)){ c->msa=0; for(int i=0;in_layers;i++) c->idx_type[i]=0; fprintf(stderr,"[MSA] indexer weights absent — full causal attention " "(add them with: convert --arch m3 --indexer for block-sparse long context)\n"); } } + /* COLI_MSA=0: force full causal attention even when the indexer weights ARE present. + * A/B switch for measuring MSA's long-context effect; mirrors the auto-detect fallback + * above (same two lines), so it rides the already-tested full-attention path. */ + if(c->msa && getenv("COLI_MSA") && atoi(getenv("COLI_MSA"))==0){ + c->msa=0; for(int i=0;in_layers;i++) c->idx_type[i]=0; + fprintf(stderr,"[MSA] COLI_MSA=0 — full causal attention (MSA disabled for A/B)\n"); + } m->kv=calloc(1,sizeof(KVState)); m->kv_start=m->kv->kv_start=calloc(NR,sizeof(int)); for(int i=layer_begin;i=0.35){ - fprintf(f,"[PROF] verdict: attention-bound (%.0f%%) — context length is the cost (DSA %s). A lower CTX helps if the workload allows.\n", - 100*f_attn, m->has_dsa?"on":"not available for this model"); + { int is_m3 = (m->c.arch==ARCH_M3); + const char *sp_name = is_m3 ? "MSA" : "DSA"; + const char *sp_on = is_m3 ? (m->c.msa?"on":"not available for this model") + : (m->has_dsa?"on":"not available for this model"); + fprintf(f,"[PROF] verdict: attention-bound (%.0f%%) — context length is the cost (%s %s). A lower CTX helps if the workload allows.\n", + 100*f_attn, sp_name, sp_on); } } else { fprintf(f,"[PROF] verdict: balanced — no phase dominates (I/O %.0f%%, matmul %.0f%%, attention %.0f%%); this config is a reasonable fit for this machine.\n", 100*f_io,100*f_emm,100*f_attn); @@ -11425,7 +11436,18 @@ int main(int argc, char **argv){ * the stored weight format: a fmt=4 grouped-int4 container still computes at * 8-bit here. Label it as compute so the banner is not misread as a storage * claim (#1183). */ - printf("== GLM C engine (glm_moe_dsa), cache=%d experts/layer | compute experts@%d-bit dense@%d-bit | idot: " IDOT_KERNEL " ==\n", cap, ebits, dbits); + /* Name the actual architecture. The engine identity prints before model_init, so + * the arch is not known yet -- peek at config.json's model_type (cheap, safe). */ + { const char *engine_name = "GLM C engine (glm_moe_dsa)"; + char cf[2100]; snprintf(cf,sizeof(cf),"%s/config.json",snap); + FILE *cfh=fopen(cf,"rb"); + if(cfh){ fseek(cfh,0,SEEK_END); long cn=ftell(cfh); fseek(cfh,0,SEEK_SET); + char *cb=(cn>0 && cn<=(1<<20))?malloc((size_t)cn+1):NULL; + if(cb){ size_t cc=fread(cb,1,(size_t)cn,cfh); cb[cc]=0; + if(strstr(cb,"minimax")) engine_name="MiniMax-M3 C engine (minimax_m3 MSA)"; + free(cb); } + fclose(cfh); } + printf("== %s, cache=%d experts/layer | compute experts@%d-bit dense@%d-bit | idot: " IDOT_KERNEL " ==\n", engine_name, cap, ebits, dbits); } g_mem_avail_boot = mem_available_gb(); #if !defined(_WIN32) if(getenv("CLUSTER_WORKERS") && *getenv("CLUSTER_WORKERS")){ diff --git a/c/family_registry.py b/c/family_registry.py index 951a5323e..9db39d7fc 100644 --- a/c/family_registry.py +++ b/c/family_registry.py @@ -653,7 +653,7 @@ def _inkling_expert_inventory(name, size, config): # resolves itself -- the legacy 8 belongs to the sister engines. # One KV slot: batched serve is not validated for M3 in this PR. limits=FamilyLimits(4096, 1048576, 1024, 16384, 1, 0, "CTX"), - capabilities=FamilyCapabilities(False, False, False, True), + capabilities=FamilyCapabilities(True, False, False, True), has_gateway_adapter=True, has_cli_adapter=True, tune_prompt_template="]~!b[]~b]user\n{prompt}[e~[\n]~b]ai\n", diff --git a/c/openai_server.py b/c/openai_server.py index 1ec8acbd9..a0f164f29 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -529,9 +529,160 @@ def parse_arch_tool_calls(reply, tools, tool_reply=None): _sideband_text, calls = parse_k3_tool_calls(tool_reply, tools) return reply.strip(), calls return parse_k3_tool_calls(reply, tools) # compatibility with pre-#1147 engines + if ARCH == "minimax_m3": + return parse_tool_calls_m3(reply, tools) return parse_tool_calls(reply, tools) +# ---- MiniMax-M3 tool calling (chat_template.jinja) ---------------------------------------- +# MiniMax-M3 uses a namespaced XML dialect, not GLM's {name}… form: +# ]<]minimax[>[ +# ]<]minimax[>[]<]minimax[>[v]<]minimax[>[]<]minimax[>[ +# ]<]minimax[>[ +# Tool results render as ]~b]tool\n…. Tools are declared in the developer +# message as {json}…. Rendered + parsed here so the MiniMax-M3 +# template can serve tools end-to-end (previously the dispatcher returned 400 for `tools`). +MM_NS = "]<]minimax[>[" +MM_TC_BEGIN = MM_NS + "" +MM_TC_END = MM_NS + "" +_MM_TC_BLOCK_RE = re.compile(re.escape(MM_TC_BEGIN) + r"\s*(.*?)\s*" + re.escape(MM_TC_END), re.DOTALL) +_MM_INVOKE_RE = re.compile(re.escape(MM_NS) + r'(.*?)' + re.escape(MM_NS) + r"", re.DOTALL) + + +def _m3_to_xml(value): + """Recursive MiniMax-M3 tool-argument XML (mirrors chat_template.jinja `to_xml`).""" + if isinstance(value, dict): + return "\n".join(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}" + for k, v in value.items() if v is not None) + if isinstance(value, (list, tuple)): + return "\n".join(f"{MM_NS}{_m3_to_xml(i)}{MM_NS}" for i in value) + if isinstance(value, bool): + return json.dumps(value) + return "" if value is None else str(value) + + +def _m3_tool_call_block(tool_calls): + """Render an assistant message's tool_calls in MiniMax-M3 invoke syntax.""" + parts = [MM_TC_BEGIN + "\n"] + for tc in (tool_calls or []): + fn = tc.get("function", tc) if isinstance(tc, dict) else {} + name = fn.get("name") or "" + parts.append(f'{MM_NS}\n') + args = fn.get("arguments", "{}") + if isinstance(args, str): + try: + args = json.loads(args) + except (json.JSONDecodeError, TypeError): + args = {} + for k, v in (args or {}).items(): + if v is None: + continue + parts.append(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}\n") + parts.append(MM_NS + "\n") + parts.append(MM_TC_END) + return "".join(parts) + + +def _m3_tools_block(tools, tool_choice=None): + """Render the MiniMax-M3 developer-side tool declaration block.""" + forced = None + if isinstance(tool_choice, dict): + forced = ((tool_choice.get("function") or {}).get("name") or tool_choice.get("name")) + tools = [t for t in (tools or []) + if ((t.get("function", t) if isinstance(t, dict) else {}).get("name") == forced)] + out = ["\n\n# Tools\nYou may call one or more tools to assist with the user query.\n" + "Here are the tools available in JSONSchema format:\n\n\n"] + for tool in (tools or []): + fn = tool.get("function", tool) if isinstance(tool, dict) else {} + clean = {k: v for k, v in fn.items() if k not in ("defer_loading", "strict")} + out.append("" + json.dumps(clean, ensure_ascii=False) + "\n") + out.append("\n\nTo call tools, wrap all invocations in a single " + + MM_TC_BEGIN + MM_TC_END + " block. Parameter values containing nested " + "objects or arrays are recursively expanded into XML elements. Example:\n\n" + + MM_TC_BEGIN + "\n" + + MM_NS + '' + MM_NS + "value-1" + MM_NS + "\n" + + MM_NS + "\n" + MM_TC_END) + if forced: + out.append(f"\n\nYou must call the function `{forced}`. Do not answer directly.") + elif tool_choice == "required": + out.append("\n\nYou must call one of the functions above. Do not answer directly.") + return "".join(out) + + +def _m3_children(text): + """Split a run of MM_NS…MM_NS blocks into (tag, inner) pairs (depth 0).""" + out = [] + i, n = 0, len(text) + ns = len(MM_NS) + while i < n: + if text.startswith(MM_NS + "<", i): + j = text.find(">", i + ns) + if j < 0: + break + tag = text[i + ns + 1:j] + open_tag, close_tag = MM_NS + "<" + tag + ">", MM_NS + "" + start, depth, k = j + 1, 1, j + 1 + while k < n: + if text.startswith(open_tag, k): + depth += 1 + k += len(open_tag) + elif text.startswith(close_tag, k): + depth -= 1 + if depth == 0: + out.append((tag, text[start:k])) + i = k + len(close_tag) + break + k += len(close_tag) + else: + k += 1 + else: + break + else: + i += 1 + return out + + +def _m3_parse_value(text): + """Parse a MiniMax-M3 XML value back into Python (mirror of _m3_to_xml).""" + text = text.strip() + if not text: + return "" + if text.startswith(MM_NS + "<"): + children = _m3_children(text) + if not children: + return text + if children[0][0] == "item": + return [_m3_parse_value(inner) for _, inner in children] + return {tag: _m3_parse_value(inner) for tag, inner in children} + return _coerce_arg(text, None) + + +def _m3_parse_args(body, types): + """Parse the top-level … argument pairs of one invoke, type-coerced.""" + args = {} + for key, raw in _m3_children(body): + raw = raw.strip() + args[key] = (_coerce_arg(raw, types.get(key)) if (MM_NS + "<") not in raw + else _m3_parse_value(raw)) + return args + + +def parse_tool_calls_m3(reply, tools=None): + """Return (content, tool_calls) for the MiniMax-M3 invoke dialect.""" + param_types = _tool_param_types(tools) + calls = [] + for block in _MM_TC_BLOCK_RE.finditer(reply): + for inv in _MM_INVOKE_RE.finditer(block.group(1)): + name = inv.group(1) + args = _m3_parse_args(inv.group(2), param_types.get(name, {})) + calls.append({"id": "call_" + uuid.uuid4().hex[:24], "type": "function", + "function": {"name": name, + "arguments": json.dumps(args, ensure_ascii=False)}}) + text = _MM_TC_BLOCK_RE.sub("", reply) + return text.strip(), calls + + + def _tool_stream_markers(): """Marker(s) that open a model tool-call block, in match order (arch-specific).""" if ARCH == "deepseek_v4": @@ -1277,14 +1428,17 @@ def render_chat_inkling(messages, enable_thinking=False, reasoning_effort=None, def render_chat_m3(messages, enable_thinking=False, reasoning_effort=None, tools=None, tool_choice=None): """Text-only subset of the MiniMax-M3 chat template (chat_template.jinja): - ]~!b[ once, then ]~b]\n[e~[\n blocks. A client "system" message + ]~!b[ once, then ]~b]\\n[e~[\\n blocks. A client "system" message maps to the `developer` role (the official template reserves `system` for the auto-injected model-identity block, which we do not fabricate here). History assistant turns carry the prefix; the open ai turn does too unless - thinking is enabled. Tool calls: not yet rendered for this family.""" - if tools: - raise APIError(400, "tools are not yet supported for the MiniMax-M3 template.", "tools") + thinking is enabled. Tool calls use the ]<]minimax[>[ XML + dialect; tool results come back as ]~b]tool\\n….""" + if tool_choice == "none": + tools = None prompt = ["]~!b["] + dev_done = False + prev_tool = False for index, message in enumerate(messages): if not isinstance(message, dict): raise APIError(400, "Each message must be an object.", f"messages.{index}") @@ -1292,14 +1446,29 @@ def render_chat_m3(messages, enable_thinking=False, reasoning_effort=None, tools raw = message.get("content") text = content_text(raw, f"messages.{index}.content") if raw is not None else "" if role in ("system", "developer"): - prompt.append(f"]~b]developer\n{text}[e~[\n") + block = f"]~b]developer\n{text}" + if tools and not dev_done: + block += _m3_tools_block(tools, tool_choice) + dev_done = True + prompt.append(block + "[e~[\n") elif role == "user": prompt.append(f"]~b]user\n{text}[e~[\n") elif role == "assistant": - prompt.append(f"]~b]ai\n{text.strip()}[e~[\n") + tc_block = (_m3_tool_call_block(message.get("tool_calls")) + if message.get("tool_calls") else "") + prompt.append(f"]~b]ai\n{text.strip()}{tc_block}[e~[\n") + elif role == "tool": + if not prev_tool: + prompt.append("]~b]tool") + prompt.append(f"\n{text}") + if index == len(messages) - 1 or messages[index + 1].get("role") != "tool": + prompt.append("[e~[\n") else: raise APIError(400, f"Unsupported message role for MiniMax-M3: {role!r}.", f"messages.{index}.role", "unsupported_role") + prev_tool = (role == "tool") + if tools and not dev_done: + prompt.insert(1, f"]~b]developer\n{_m3_tools_block(tools, tool_choice)}[e~[\n") prompt.append("]~b]ai\n" + ("" if enable_thinking else "")) return "".join(prompt) From 04bd7c905bd5798bab4b05ce2b32832fe44dbd70 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sat, 15 Aug 2026 18:27:39 +0200 Subject: [PATCH 16/22] m3 tools: byte-exact template rendering, full example block, stream marker Three adjustments on top of the tool-call support from fork PR #1, all measured against the official chat_template.jinja rendered via jinja2: - invoke/to_xml rendering: the template's for-loops are whitespace-trimmed on both sides, so argument pairs and nested elements concatenate with NO newlines; dropped the extra \n after and between argument pairs (the parser was already layout-agnostic). - developer-block example: the template's full two-invoke example including the nested param-2 item block, byte-identical. - _tool_stream_markers: minimax holds the visible stream at the ]<]minimax[>[ opener. The arch-dispatched streaming scan landed on dev after this PR's base; without a marker a streamed reply would leak the raw tool-call block to the client token by token. Validated against the jinja reference render: tools declaration, assistant tool_call block, and tool-response run all byte-identical; parsing the template-rendered block round-trips nested objects/arrays/bools with schema type coercion. --- c/openai_server.py | 24 ++++++++++++++++++------ 1 file changed, 18 insertions(+), 6 deletions(-) diff --git a/c/openai_server.py b/c/openai_server.py index a0f164f29..9f9286a1d 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -552,10 +552,10 @@ def parse_arch_tool_calls(reply, tools, tool_reply=None): def _m3_to_xml(value): """Recursive MiniMax-M3 tool-argument XML (mirrors chat_template.jinja `to_xml`).""" if isinstance(value, dict): - return "\n".join(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}" - for k, v in value.items() if v is not None) + return "".join(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}" + for k, v in value.items() if v is not None) if isinstance(value, (list, tuple)): - return "\n".join(f"{MM_NS}{_m3_to_xml(i)}{MM_NS}" for i in value) + return "".join(f"{MM_NS}{_m3_to_xml(i)}{MM_NS}" for i in value) if isinstance(value, bool): return json.dumps(value) return "" if value is None else str(value) @@ -567,7 +567,7 @@ def _m3_tool_call_block(tool_calls): for tc in (tool_calls or []): fn = tc.get("function", tc) if isinstance(tc, dict) else {} name = fn.get("name") or "" - parts.append(f'{MM_NS}\n') + parts.append(f'{MM_NS}') args = fn.get("arguments", "{}") if isinstance(args, str): try: @@ -577,7 +577,7 @@ def _m3_tool_call_block(tool_calls): for k, v in (args or {}).items(): if v is None: continue - parts.append(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}\n") + parts.append(f"{MM_NS}<{k}>{_m3_to_xml(v)}{MM_NS}") parts.append(MM_NS + "\n") parts.append(MM_TC_END) return "".join(parts) @@ -600,7 +600,17 @@ def _m3_tools_block(tools, tool_choice=None): + MM_TC_BEGIN + MM_TC_END + " block. Parameter values containing nested " "objects or arrays are recursively expanded into XML elements. Example:\n\n" + MM_TC_BEGIN + "\n" - + MM_NS + '' + MM_NS + "value-1" + MM_NS + "\n" + + MM_NS + '' + + MM_NS + "value-1" + MM_NS + "" + + MM_NS + "" + + MM_NS + "" + + MM_NS + "val-a" + MM_NS + "" + + MM_NS + "val-b" + MM_NS + "" + + MM_NS + "" + + MM_NS + "" + + MM_NS + "\n" + + MM_NS + '' + + MM_NS + "value-1" + MM_NS + "" + MM_NS + "\n" + MM_TC_END) if forced: out.append(f"\n\nYou must call the function `{forced}`. Do not answer directly.") @@ -689,6 +699,8 @@ def _tool_stream_markers(): return ("<" + DSV4_DSML + "tool_calls", "<" + DSV4_DSML + "invoke") if ARCH == "kimi": return (K3_TOOLS_OPEN,) + if ARCH == "minimax_m3": + return (MM_TC_BEGIN,) return (BOX_START,) From 29e68aaf00ca8c548e88bd7778dfdf7c3d29356b Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 16 Aug 2026 19:54:08 +0200 Subject: [PATCH 17/22] fix(m3): resolve EOS in one helper for run_text/run_serve/run_serve_mux The [e~[ fallback lived in run_serve only, but openai_server launches the engine with SERVE_BATCH=1 -> run_serve_mux, and coli run goes through run_text: both kept eos=-1 on an M3 tokenizer. With -1 the #401 stop filter in sample.h is skipped and every tokenizer special token is armed as a hard stop, so the first tool-call marker ends the turn. One tok_eos_resolve() used by all three entry points. (Review B1.) --- c/colibri.c | 18 ++++++++++++++---- 1 file changed, 14 insertions(+), 4 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 6d1d7d1f2..9dd42f5c0 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -8443,12 +8443,23 @@ static void run_replay(Model *m, const int *full, int nfull, int np){ #endif } +/* EOS for EVERY generation entry point (run_text / run_serve / run_serve_mux): + * GLM's <|endoftext|> first, then MiniMax-M3's [e~[. One helper on purpose — a -1 + * EOS is not benign downstream: sample.h's batched-serve stop filter (#401) is + * gated on tok_eos>=0, so an unresolved EOS arms EVERY tokenizer special token as + * a hard stop and the first M3 tool-call marker would end the turn. */ +static int tok_eos_resolve(Tok *T){ + int eos=tok_id_of(T,"<|endoftext|>"); + if(eos<0) eos=tok_id_of(T,"[e~["); /* MiniMax-M3 end-of-sequence */ + return eos; +} + /* generazione reale: tokenizza PROMPT, prefill + decode greedy con stop su EOS, * detokenizza e stampa il testo in streaming. */ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){ Cfg *c=&m->c; char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); Tok T; tok_load(&T,tkp); - int eos=tok_id_of(&T,"<|endoftext|>"); + int eos=tok_eos_resolve(&T); stops_arm_tok(&m->c, eos, &T); grammar_setup(&g_grd,&T); /* metodo F: GRAMMAR=file.gbnf (#48) */ if(g_temp<0) g_temp=0.7f; /* auto: 0.7, NON l'1.0 ufficiale — la coda della @@ -9228,7 +9239,7 @@ static int mux_submit(Model *m, Tok *T, ServeCtx *ctx, ServeReq *req, GrDraft *g static void run_serve_mux(Model *m, const char *snap){ char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); - Tok T; tok_load(&T,tkp); int eos=tok_id_of(&T,"<|endoftext|>"); stops_arm_tok(&m->c,eos,&T); + Tok T; tok_load(&T,tkp); int eos=tok_eos_resolve(&T); stops_arm_tok(&m->c,eos,&T); int maxctx=getenv("CTX")?atoi(getenv("CTX")):4096; int nctx=getenv("KV_SLOTS")?atoi(getenv("KV_SLOTS")):1; if(nctx<1||nctx>512){fprintf(stderr,"KV_SLOTS must be between 1 and 512\n");exit(2);} @@ -9412,8 +9423,7 @@ static void run_serve(Model *m, const char *snap){ double t_serve0=now_s(); /* PROF: wall base for the exit-time profile_print */ char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap); Tok T; tok_load(&T,tkp); - int eos=tok_id_of(&T,"<|endoftext|>"); - if(eos<0) eos=tok_id_of(&T,"[e~["); /* MiniMax-M3 end-of-sequence */ + int eos=tok_eos_resolve(&T); stops_arm_tok(&m->c, eos, &T); grammar_setup(&g_grd,&T); /* metodo F: GRAMMAR=file.gbnf (#48) */ if(g_temp<0) g_temp=0.7f; /* auto: 0.7, NON l'1.0 ufficiale — la coda della From 034d7bce805491a1173c71bdac6dc9e089c9c95b Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 16 Aug 2026 19:58:56 +0200 Subject: [PATCH 18/22] fix(plan): size the GQA KV + MSA index cache from the registry geometry build_plan sized KV from the MLA keys only (kv_lora_rank/qk_rope_head_dim/ qk_nope_head_dim/v_head_dim), so an M3 container planned with kv_bytes=0, kv_buffer=0 and configured_experts=0: plan/doctor/--auto-tier under- reserved ~1.2 GB at 60L x 4096 ctx and the expert-cap clamp never fired. Since #1103 those formulas belong to the family descriptor, so this lands in _minimax_geometry rather than as a branch in build_plan -- the planner itself is untouched. The geometry mirrors the engine: Lc/Rc hold full K and V rows (num_key_value_heads*head_dim each) over n_layers+1 rows, the MSA index-key cache exists only on the sparse layers and without the MTP row (kv_alloc allocates Ic over n_layers and skips !idx_type), and the sparse set is derived the way load_cfg derives it -- from moe_layer_freq, because M3's sparse attention layers ARE its MoE layers. Workspace is attention_gqa's scratch set plus the indexer's own projections. Two test layers: the registry test pins the arithmetic (and now asserts against the shipped descriptor instead of the double the registry landed with), and test_resource_plan runs a real M3 container end to end, so it also fails if build_plan ever stops routing M3 through the geometry. MLA arithmetic is unchanged. (Review B2.) --- c/family_registry.py | 27 ++++++++++++- c/tests/test_family_registry.py | 14 +++++-- c/tests/test_resource_plan.py | 69 +++++++++++++++++++++++++++++++++ 3 files changed, 106 insertions(+), 4 deletions(-) diff --git a/c/family_registry.py b/c/family_registry.py index 9db39d7fc..e06c2c9cd 100644 --- a/c/family_registry.py +++ b/c/family_registry.py @@ -405,11 +405,23 @@ def _minimax_geometry(config, context, _model_dir): state = (layers + 1) * context * 2 * kv_heads * head_dim * 4 + The MSA Lightning Indexer adds one index-key row per token, but only on + the sparse layers and without the MTP row (kv_alloc allocates Ic over + n_layers and skips !idx_type layers). The engine derives sparsity from + moe_layer_freq -- M3's sparse attention layers ARE its MoE layers, so + idx_type[i] = (i >= first_dense) where first_dense is the index of the + first non-zero entry, else 0 (colibri.c load_cfg, ARCH_M3 branch): + + state += (layers - first_dense) * context * sparse_index_dim * 4 + Workspace mirrors attention_gqa()'s scratch set with S == context at prefill: q and ctx are heads-wide, k and v are kv_heads-wide, plus the - per-row score buffer (heads * window, window <= context). + per-row score buffer (heads * window, window <= context). On sparse + layers the indexer adds its own projections, index_heads * index_dim + wide for the query side and index_dim for the key side. ws = context * ((2 * heads + 2 * kv_heads) * head_dim + heads) * 4 + + context * (index_heads + 1) * sparse_index_dim * 4 Experts: configured_experts = num_local_experts (M3's spelling of n_routed_experts). @@ -422,6 +434,19 @@ def _minimax_geometry(config, context, _model_dir): state = (layers + 1) * context * 2 * kv_heads * head_dim * 4 workspace = context * ((2 * heads + 2 * kv_heads) * head_dim + heads) * 4 + + sparse = config.get("sparse_attention_config") + if isinstance(sparse, dict) and sparse.get("use_sparse_attention"): + index_dim = _required_int(sparse, "sparse_index_dim", "minimax_m3") + index_heads = _optional_int(sparse, "sparse_num_index_heads", 1, 1) + freq = config.get("moe_layer_freq") + first_dense = 0 + if isinstance(freq, list): + first_dense = next((i for i, value in enumerate(freq) if value), len(freq)) + first_dense = min(first_dense, layers) + state += (layers - first_dense) * context * index_dim * 4 + workspace += context * (index_heads + 1) * index_dim * 4 + return PlannerGeometry(state, 0, workspace, experts) diff --git a/c/tests/test_family_registry.py b/c/tests/test_family_registry.py index 7d763068b..0babbc5e6 100644 --- a/c/tests/test_family_registry.py +++ b/c/tests/test_family_registry.py @@ -129,6 +129,11 @@ def test_minimax_shares_colibri_without_becoming_glm(self): "num_local_experts": 4, "num_experts_per_tok": 2, "moe_layer_freq": [0, 1], + "sparse_attention_config": { + "use_sparse_attention": True, + "sparse_index_dim": 8, + "sparse_attention_freq": [0, 1], + }, } # minimax_m3 is a registered family now, so the fixture this test # landed with would collide on its id. Assert against the production @@ -144,9 +149,12 @@ def test_minimax_shares_colibri_without_becoming_glm(self): "model_dir": "."})() geometry = planner_geometry(resolved, 32) self.assertEqual(geometry.configured_experts, 4) - # 3 rows (2 layers + the MTP row) x 32 tokens x (K + V) x 2 kv heads - # x 8 head_dim x 4 bytes. No latent compression: GQA caches full rows. - self.assertEqual(geometry.context_state_bytes, 12_288) + # 12_288: 3 rows (2 layers + the MTP row) x 32 tokens x (K + V) x 2 kv + # heads x 8 head_dim x 4 bytes -- no latent compression, GQA caches + # full rows. Plus 1_024 for the MSA index keys, which exist on the one + # sparse layer only (moe_layer_freq [0, 1] -> first_dense 1) and get no + # MTP row: 1 x 32 x 8 x 4. + self.assertEqual(geometry.context_state_bytes, 13_312) def test_olmoe_fixture_models_conventional_fp32_kv_cache(self): # OLMoE keeps a full K and V cache per layer, sized at num_attention_heads diff --git a/c/tests/test_resource_plan.py b/c/tests/test_resource_plan.py index 347ee1f57..85f6029c7 100644 --- a/c/tests/test_resource_plan.py +++ b/c/tests/test_resource_plan.py @@ -875,5 +875,74 @@ def sysctl(command, **kwargs): self.assertEqual(physical_cpu_count(), 8) +class M3ResourcePlanTest(unittest.TestCase): + """#601 review B2: the planner sized KV from the MLA keys only, so a GQA + (MiniMax-M3) container planned with kv_bytes=0 and configured_experts=0 -- + the RAM reservation ran ~1 GB short at 60L x 4096 ctx and the expert-cap + clamp never fired. The formulas now live in the registry's + _minimax_geometry; these run a real container end to end, so they fail if + build_plan stops routing M3 through it -- not only if the geometry is + wrong.""" + + def setUp(self): + self.tmp = tempfile.TemporaryDirectory() + self.model = Path(self.tmp.name) + (self.model / "config.json").write_text(json.dumps({ + "model_type": "minimax_m3", + "num_hidden_layers": 2, + "num_attention_heads": 4, + "num_key_value_heads": 2, + "head_dim": 8, + "num_local_experts": 3, + "moe_layer_freq": [0, 1], + "sparse_attention_config": { + "use_sparse_attention": True, "sparse_index_dim": 4, + "sparse_num_index_heads": 2}, + })) + write_shard(self.model / "model.safetensors", [ + ("model.embed_tokens.weight", 100), + ("model.layers.0.self_attn.q_proj.weight", 200), + ("model.layers.1.mlp.experts.0.gate_proj.weight", 30), + ("model.layers.1.mlp.experts.1.gate_proj.weight", 30), + ("model.layers.1.mlp.experts.2.gate_proj.weight", 30), + ]) + + def tearDown(self): + self.tmp.cleanup() + + def test_gqa_kv_is_reserved(self): + context = 32 + plan = build_plan(self.model, ram_gb=16, context=context, + available_memory=32 * GB, available_disk=100 * GB, + gpus=[], physical_cpus=8, cpu_sockets=1) + # K and V rows on 2 layers + the MTP row kv_alloc also keeps, then the + # MSA index keys on the one sparse layer (moe_layer_freq [0, 1]), which + # gets no MTP row. + kv_bytes = (2 + 1) * context * 2 * 2 * 8 * 4 + 1 * context * 4 * 4 + # attention_gqa scratch: q + ctx (heads-wide), k + v (kv-heads-wide), + # the per-row scores, and the indexer's own q/k projections. + kv_buffer = (context * ((2 * 4 + 2 * 2) * 8 + 4) * 4 + + context * (2 + 1) * 4 * 4) + ram = plan["tiers"]["ram"] + self.assertEqual(ram["sequence_state_bytes"], kv_bytes) + self.assertEqual(ram["workspace_bytes"], kv_buffer) + max_expert = plan["model"]["max_expert_bytes"] + self.assertEqual(ram["runtime_bytes"], + int(1.2 * GB + 2.5 * GB + 64 * max_expert + + kv_bytes + kv_buffer)) + + def test_container_resolves_to_the_minimax_family(self): + plan = build_plan(self.model, ram_gb=16, context=32, + available_memory=32 * GB, available_disk=100 * GB, + gpus=[], physical_cpus=8, cpu_sockets=1) + self.assertEqual(plan["model"]["family_id"], "minimax_m3") + + def test_expert_cap_clamps_to_num_local_experts(self): + plan = build_plan(self.model, ram_gb=16, context=32, + available_memory=32 * GB, available_disk=100 * GB, + gpus=[], physical_cpus=8, cpu_sockets=1) + self.assertEqual(plan["tiers"]["ram"]["cache_slots_per_layer"], 3) + + if __name__ == "__main__": unittest.main() From bf80ed814a38a3723e427f4b2ed24b665ffefb4a Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 16 Aug 2026 20:01:28 +0200 Subject: [PATCH 19/22] fix(convert): keep pre-M3 outdirs resumable, re-anchor layer_idx Two silent behavior changes for people not using MiniMax leave with this: - params manifest: bits_map["idx"] was always recorded (default 8) and "arch" always written, so check_or_record_params refused to resume any conversion started with the pre-M3 converter. The idx int8 default now lives at the consumer in convert_shard and both keys are recorded only when non-default: a default GLM/DeepSeek run builds a byte-identical params dict, while a REAL flag change (or --arch m3 on a GLM outdir) still aborts. "idx" classification itself only ever fires on M3's self_attn.index_* names (GLM's DSA indexer is "indexer", skipped or converted via --indexer), so no other arch's output bytes change. - layer_idx() matched "layers." anywhere; back to anchored "model.layers.N" plus the one vetted raw form, MiniMax-VL's "language_model.model.layers.N". (Review B3.) --- c/tools/convert_fp8_to_int4.py | 30 ++++++++++++++++++++---------- 1 file changed, 20 insertions(+), 10 deletions(-) diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 4e4f1a099..3c726d77e 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -189,13 +189,14 @@ def quant_int2(w, bits): # -> (qbytes U8 [O*ceil(I/4)], s # ---------- classificazione dei tensori ---------- def layer_idx(name): - # find "layers." anywhere: handles both "model.layers.N..." (mapped) and the raw - # "language_model.model.layers.N..." (MiniMax-M3 VL prefix) forms. + # ANCHORED match: "model.layers.N..." (mapped) or the raw MiniMax-M3 VL prefix + # "language_model.model.layers.N..." -- NOT "layers." anywhere, which would + # change classification for every arch on names this converter never vetted. p = name.split(".") - for i in range(len(p) - 1): - if p[i] == "layers": - try: return int(p[i + 1]) - except ValueError: return -1 + if len(p) > 1 and p[0] == "language_model" and p[1] == "model": p = p[1:] + if len(p) > 2 and p[0] == "model" and p[1] == "layers": + try: return int(p[2]) + except ValueError: return -1 return -1 def classify(name, n_layers, keep_mtp=False, keep_idx=False): @@ -399,10 +400,14 @@ def convert_shard(path, out_dict, n_layers, ebits, io_bits, xbits, # otherwise fall back to the classic ebits/xbits/io_bits scheme. if bits_map and kind in bits_map: bits = bits_map[kind] + elif kind == "idx": + # MSA scoring indexer (M3 self_attn.index_*): block selection is + # DISCRETE top-k, so noise flips choices -- int8 default, not ebits. + bits = 8 else: bits = io_bits if kind == "io" else xbits if kind == "x" else ebits # Any unknown kind that fell through classify as "q" - if bits_map and kind not in bits_map and kind not in ("io", "x", "sh", "o", "kvb", "attn", "dmlp"): + if bits_map and kind not in bits_map and kind not in ("io", "x", "sh", "o", "kvb", "attn", "dmlp", "idx"): bits = ebits # Per-projection override for routed experts, applied on top of the type-level bits. if kind == "x" and PROJ_BITS: # e.g. up_proj -> 3 (int3-g64) while gate/down stay 4 @@ -607,7 +612,10 @@ def main(): if a.kvb_bits is not None: bits_map["kvb"] = a.kvb_bits if a.attn_bits is not None: bits_map["attn"] = a.attn_bits if a.dmlp_bits is not None: bits_map["dmlp"] = a.dmlp_bits - bits_map["idx"] = a.idx_bits # always set: default 8 (see --idx-bits) + # "idx" is recorded only when overridden: the int8 default lives at the consumer + # in convert_shard, so default runs keep the params manifest byte-identical to + # pre-M3 converters and in-progress GLM/DeepSeek outdirs stay resumable. + if a.idx_bits != 8: bits_map["idx"] = a.idx_bits if bits_map: print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items()))) @@ -772,7 +780,8 @@ def get_slice(s, n): return None # 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, - "proj_bits": dict(PROJ_BITS), "arch": a.arch} + "proj_bits": dict(PROJ_BITS)} + if a.arch != "glm": params["arch"] = a.arch # default omitted: pre-M3 manifests stay resumable prog_path = os.path.join(a.outdir, f".{prefix}progress.json") prog = {} if os.path.exists(prog_path): @@ -1137,7 +1146,8 @@ def _download_single(url, fn, out, part, expected): shutil.rmtree(tmp, ignore_errors=True); print("[IDX] DONE."); return 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, - "proj_bits": dict(PROJ_BITS), "arch": a.arch} + "proj_bits": dict(PROJ_BITS)} + if a.arch != "glm": params["arch"] = a.arch # default omitted: pre-M3 manifests stay resumable if not check_or_record_params(a.outdir, "out-", params): return for i, sh in enumerate(shards): if free_gb(a.outdir) < a.min_free_gb: From 7e3f4998f1bdea44a6bfe005b94b909ecb3185d5 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 16 Aug 2026 20:09:32 +0200 Subject: [PATCH 20/22] feat(m3): CI oracle gate, and the shared-binary gates the registry cannot infer M3 rides the colibri binary, so a GLM-side refactor of the shared hot paths (rmsnorm/act_glu/expert_ffn/MSA) could break it with every GLM test green. Now something notices (Review B4): - make m3-tiny-generate / m3-tiny-check: tiny random checkpoint -> --arch m3 int8 conversion -> numpy oracle -> the engine's REF/TF gate; tests/test_m3_tiny.py enforces token-exact counts (the engine itself always exits 0). New ci.yml job runs it on Linux with pinned deps (tools/requirements-m3-tiny.txt). Everything the earlier revision of this commit added to coli and openai_server -- banner roster, model_arch, engine dict, build target, tune prompt, model ids, --arch choices, the KV-slot guard, the cap sentinel -- is now supplied by the FamilyDescriptor and is gone from this diff. What the descriptor could NOT express was the handful of launcher gates written as `arch == "glm"`, which is a claim about the engine binary, not about the family. They now ask the registry the question they actually mean, via colibri_core(): family.engine_group == family_by_id("glm").engine_group engine_group is the registry's own statement that two families share an engine implementation -- and therefore its env contract, its CAP channel and its byte-protocol REPL. Four sites move: env_for_engine (M3 was about to get the sister-engine OMP treatment for a binary that sizes its own team), operator_cap (CAP= was silently ignored), cmd_chat (M3 would have spawned a gateway server instead of the byte-protocol REPL), and engine_for's COLI_DOCKER_GLM_ONLY guard, which told users a GLM-only image cannot run M3 -- it can; it is the same binary. Two gates deliberately keep the narrow `arch == "glm"` test: `coli tune`'s replay protocol (M3 tunes through the persistent-rotation path; giving it GLM's fixed replay means touching six more branches in autotune.py, unvalidated here) and need_worker_model (cluster expert workers are not validated for M3, and refusing says so). test_launcher_dispatch needs no M3 entry now: the registry asserts each family resolves to its own declared engine_artifact, and the intentional share is declared by the descriptor rather than by a set in the test. --- .github/workflows/ci.yml | 21 +++++++++++ c/Makefile | 16 ++++++++ c/coli | 20 ++++++++-- c/openai_server.py | 2 +- c/tests/test_m3_tiny.py | 63 ++++++++++++++++++++++++++++++++ c/tools/requirements-m3-tiny.txt | 8 ++++ 6 files changed, 125 insertions(+), 5 deletions(-) create mode 100644 c/tests/test_m3_tiny.py create mode 100644 c/tools/requirements-m3-tiny.txt diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 2f1e5c54b..41d64d6df 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -17,6 +17,27 @@ jobs: - name: C test suite run: cd c && make test-c + # MiniMax-M3 rides the colibri binary (ARCH_M3 auto-detect), so a GLM-side + # refactor of the shared hot paths (rmsnorm/act_glu/expert_ffn/MSA) can break + # M3 while every GLM test stays green. This gate is the only thing that + # notices: tiny random M3 checkpoint -> converter -> numpy oracle -> the + # engine's teacher-forcing gate, token-exact on every position (#601). + m3-tiny: + name: MiniMax-M3 (generated tiny oracle) + runs-on: ubuntu-latest + timeout-minutes: 15 + steps: + - uses: actions/checkout@v4 + - uses: actions/setup-python@v5 + with: + python-version: '3.13' + cache: pip + cache-dependency-path: c/tools/requirements-m3-tiny.txt + - name: Install pinned tiny-fixture dependencies + run: python -m pip install -r c/tools/requirements-m3-tiny.txt + - name: Generate fixture and run the M3 oracle gate + run: make -C c m3-tiny-check + # Every engine, on every platform a release archive is built for. Before this, # CI built colibri + inkling on Linux only and kimi_k3 on nothing -- so the fact # that inkling and kimi_k3 did not compile under msys2/UCRT64 at all was invisible diff --git a/c/Makefile b/c/Makefile index 06fd6882c..ee0a62f5a 100644 --- a/c/Makefile +++ b/c/Makefile @@ -668,6 +668,22 @@ deepseek-v4-tiny-generate: @echo "SKIP deepseek-v4-tiny-generate: V4 runtime requires x86-64/aarch64 Linux or Windows/MSYS2" endif +# MiniMax-M3 CI gate (#601): tiny random checkpoint -> int8 conversion -> numpy +# oracle -> the engine's REF/TF gate, token-exact on every prefill and decode +# position (the driver enforces the counts; the engine itself always exits 0). +# Deps: tools/requirements-m3-tiny.txt (torch generates, numpy+safetensors +# convert and run the oracle). M3 rides the main colibri binary, so this needs +# no engine platform gate beyond colibri's own. +.PHONY: m3-tiny-generate m3-tiny-check +m3-tiny-generate: + rm -rf m3tiny m3tiny_i8 + $(PYTHON) tools/make_m3tiny.py m3tiny + $(PYTHON) tools/convert_fp8_to_int4.py --arch m3 --indir m3tiny --outdir m3tiny_i8 --ebits 8 --io-bits 8 + $(PYTHON) tools/oracle_m3.py m3tiny_i8 + +m3-tiny-check: m3-tiny-generate colibri$(EXE) + $(PYTHON) tests/test_m3_tiny.py --binary ./colibri$(EXE) --snap m3tiny_i8 --ref ref_m3.json + # Convenience aliases — 'glm' kept for backward compatibility. # On POSIX EXE is empty, so `colibri: colibri$(EXE)` becomes a self-dependency # and marking it phony also forces the real binary to rebuild on every invocation. diff --git a/c/coli b/c/coli index e369ee4ea..bf8fb9bd6 100755 --- a/c/coli +++ b/c/coli @@ -313,10 +313,22 @@ def redraw_prompt_box(message, width): def model_arch(model): return resolve_model(model).descriptor.id +def colibri_core(family): + """Is this family served by the colibri binary itself, rather than by a + sister engine? + + engine_group is the registry's statement that two families share an engine + implementation, so they also share that engine's env contract, its CAP + channel and its byte-protocol REPL. Keyed on the group rather than on + id == "glm" so a future third family on this binary inherits the same + handling by declaring the group, not by being added to a tuple here.""" + if isinstance(family, str): family = family_by_id(family) + return family.engine_group == family_by_id("glm").engine_group + def engine_for(model): family = resolve_model(model).descriptor if os.environ.get("COLI_ENGINE"): - if (family.id != "glm" and + if (not colibri_core(family) and os.environ.get("COLI_DOCKER_GLM_ONLY") == "1"): raise UnknownFamilyError( f"this image contains only the GLM engine; {family.display_name} " @@ -396,7 +408,7 @@ def cap_for_launch(explicit_cap, env, fallback): def operator_cap(a, arch): """Return a real operator cap, including GLM's documented CAP channel.""" cap=getattr(a,"cap",None) - if cap is not None or arch!="glm": + if cap is not None or not colibri_core(arch): return cap try: cap=int(os.environ.get("CAP","0")) @@ -405,7 +417,7 @@ def operator_cap(a, arch): return cap if cap else None def env_for_engine(a, arch, plan=None): - if arch == "glm": + if colibri_core(arch): return env_for(a) explicit_env = set(os.environ) env = os.environ.copy() @@ -1404,7 +1416,7 @@ def cmd_chat(a): arch=family.id if not family.has_gateway_adapter: sys.exit(f"{family.display_name}: gateway adapter is not wired") - if arch!="glm": + if not colibri_core(family): engine=engine_for(a.model) need_model(a.model,engine) model_id=family.default_model_id diff --git a/c/openai_server.py b/c/openai_server.py index 9f9286a1d..33c65c8b6 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -720,7 +720,7 @@ def _tool_hold(): ARCH = "glm" # set in main(): a family id from family_registry (glm | inkling | - # kimi | olmoe | qwen36 | deepseek_v4) + # kimi | olmoe | qwen36 | deepseek_v4 | minimax_m3) INK_THINK, INK_TEXT = "<|content_thinking|>", "<|content_text|>" diff --git a/c/tests/test_m3_tiny.py b/c/tests/test_m3_tiny.py new file mode 100644 index 000000000..7e2d71be5 --- /dev/null +++ b/c/tests/test_m3_tiny.py @@ -0,0 +1,63 @@ +#!/usr/bin/env python3 +"""#601 CI gate for the MiniMax-M3 port: run the engine on the converted tiny +M3 container in teacher-forcing mode and require EVERY prefill and decode +position to match the numpy oracle's argmax. + +The fixture chain (make m3-tiny-generate) is tools/make_m3tiny.py -> +tools/convert_fp8_to_int4.py --arch m3 -> tools/oracle_m3.py; this driver only +runs the engine, because the engine's REF/TF gate prints its counts but always +exits 0 -- the pass/fail decision lives here. + +IDOT=0 on purpose: int8 activation quantization can flip a borderline argmax +(~0.06 logit margin on the tiny fixture), so the gate compares the exact f32 +path -- documented in tools/oracle_m3.py. Everything under __main__ so +`unittest discover` can import this file without running the engine. +""" + +from __future__ import annotations + +import argparse +import os +import re +import subprocess +import sys + + +def main() -> int: + ap = argparse.ArgumentParser() + ap.add_argument("--binary", default="./colibri") + ap.add_argument("--snap", default="m3tiny_i8") + ap.add_argument("--ref", default="ref_m3.json") + a = ap.parse_args() + + env = dict(os.environ, SNAP=a.snap, REF=a.ref, + TF="1", TF_DECODE="1", IDOT="0") + proc = subprocess.run([a.binary, "8"], env=env, text=True, + encoding="utf-8", errors="replace", + stdout=subprocess.PIPE, stderr=subprocess.PIPE, + timeout=600) + sys.stdout.write(proc.stdout) + sys.stderr.write(proc.stderr) + + prefill = re.search(r"PREFILL \(teacher-forcing\) C vs oracle: (\d+)/(\d+)", + proc.stdout) + decode = re.search(r"DECODE \(incremental\) C vs oracle: (\d+)/(\d+)", + proc.stdout) + problems = [] + if proc.returncode: + problems.append(f"engine exited {proc.returncode}") + for label, match in (("prefill", prefill), ("decode", decode)): + if not match: + problems.append(f"{label} gate line missing from engine output") + elif int(match.group(2)) == 0 or match.group(1) != match.group(2): + problems.append(f"{label} {match.group(1)}/{match.group(2)}") + if problems: + print(f"m3-tiny-check: FAIL ({'; '.join(problems)})") + return 1 + print(f"m3-tiny-check: OK (prefill {prefill.group(0).split(': ')[1]}, " + f"decode {decode.group(1)}/{decode.group(2)})") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/c/tools/requirements-m3-tiny.txt b/c/tools/requirements-m3-tiny.txt new file mode 100644 index 000000000..3278f5900 --- /dev/null +++ b/c/tools/requirements-m3-tiny.txt @@ -0,0 +1,8 @@ +--extra-index-url https://download.pytorch.org/whl/cpu +# Same torch/safetensors pins as requirements-deepseek-v4-tiny.txt (the tiny +# M3 fixture chain is torch to generate, numpy+safetensors to convert and to +# run the oracle). numpy is pinned explicitly: torch no longer depends on it. +torch==2.13.0+cpu; sys_platform != "darwin" +torch==2.13.0; sys_platform == "darwin" +safetensors==0.8.0 +numpy==2.3.2 From 0ec2d2e65e277a49ecd41211dacb68161958c180 Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Sun, 16 Aug 2026 20:14:12 +0200 Subject: [PATCH 21/22] chore(m3): drop -Wno-missing-field-initializers, quiet VK warning, docs Review follow-ups: - The Makefile suppression existed only because the VK GQA commit widened struct PC with act/alpha/limit push constants and left the existing positional initializers short. Name every field at all 16 sites (the gate_up dispatches overwrite act/alpha/limit right after, so semantics are unchanged -- the fields were implicitly zero before) and remove the flag from all five platform CFLAGS lines. gcc 15 builds warning-free. - kv_alloc's index-cache calloc: (unsigned) cast gives value-range propagation a provable bound so -Walloc-size-larger-than stays quiet on inlined paths (the alloc is dev's own line; our wider "|| c->msa" guard changed inlining enough to surface it on the reviewer's gcc). - docs/minimax-m3.md: family doc covering the shared-binary design, GQA/ MSA/swigluoai/router notes, the CTX=4096 default and the MSA 2048-token exactness bound, tool-calling dialect, and the m3-tiny CI gate. --- c/Makefile | 10 +++---- c/backend_vulkan.c | 32 ++++++++++---------- c/colibri.c | 4 ++- docs/minimax-m3.md | 75 ++++++++++++++++++++++++++++++++++++++++++++++ 4 files changed, 99 insertions(+), 22 deletions(-) create mode 100644 docs/minimax-m3.md diff --git a/c/Makefile b/c/Makefile index ee0a62f5a..ba4b1a6d6 100644 --- a/c/Makefile +++ b/c/Makefile @@ -62,7 +62,7 @@ $(warning libomp not found: building single-threaded. For multithreading: brew i OMPC = OMPL = endif -CFLAGS = -O3 $(OMPC) -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) +CFLAGS = -O3 $(OMPC) -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) # Opt-in: ARCH=native appends -mcpu=native (arm64 clang uses -mcpu, not -march), # which unlocks the i8mm SMMLA int8/int4 dot kernels in colibri.c. ARCH unset -> # no -mcpu, default build byte-identical. Apple clang knows apple-m4 / native. @@ -96,7 +96,7 @@ else ifneq ($(IS_WIN),) # a v3 build simply compiles out the VNNI path - safe on any x86-64. CC = gcc ARCH ?= x86-64-v3 -CFLAGS = -D_FILE_OFFSET_BITS=64 -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) +CFLAGS = -D_FILE_OFFSET_BITS=64 -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) # -lpsapi: compat.h calls GetProcessMemoryInfo (rss_gb). It's linked via # #pragma comment(lib,"psapi.lib") for MSVC, but MinGW gcc ignores that pragma # (warns), so link psapi explicitly or the build fails undefined-reference on @@ -115,7 +115,7 @@ ifneq (,$(PPC64)) # (validated token-exact vs the transformers oracle on a POWER8 S824). CC = gcc ARCH ?= native -CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) +CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto @@ -168,7 +168,7 @@ ARCHFLAG := -march=$(ARCH) else ARCHFLAG := -mcpu=$(ARCH) endif -CFLAGS = -O3 $(ARCHFLAG) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) +CFLAGS = -O3 $(ARCHFLAG) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto @@ -183,7 +183,7 @@ CC = gcc # ARCH=x86-64 -> massima compatibilita' (niente AVX2: usa il path scalare di fallback). # -pthread: Linux lo tira dentro via -fopenmp, i *BSD no e le pthread_* non risolvono (#219). ARCH ?= native -CFLAGS = -O3 -march=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function -Wno-missing-field-initializers $(EXTRA_CFLAGS) +CFLAGS = -O3 -march=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function $(EXTRA_CFLAGS) LDFLAGS = -lm -fopenmp -pthread $(EXTRA_LDFLAGS) ifeq ($(LTO),1) CFLAGS += -flto diff --git a/c/backend_vulkan.c b/c/backend_vulkan.c index cd9980789..32b18fe7a 100644 --- a/c/backend_vulkan.c +++ b/c/backend_vulkan.c @@ -642,7 +642,7 @@ int coli_vk_matmul(ColiVkTensor **tensor, float *y, const float *x, VKCHECK(vkBeginCommandBuffer(G.cmd, &begin), "beginCmd"); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); - struct PC pc = {fmt, S, I, O, t->rowWords, t->gs}; + struct PC pc = {fmt, S, I, O, t->rowWords, t->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); /* Grid-stride shader: one subgroup per output row (~8 rows/workgroup at wave32). * Launch ~O/8 workgroups for occupancy; the shader loops to cover any O / wave width. */ @@ -706,7 +706,7 @@ int coli_vk_gate_up(ColiVkTensor **gate, ColiVkTensor **up, float *hidden, const VKCHECK(vkBeginCommandBuffer(G.cmd, &begin), "beginCmd"); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gu, 0, 1, &G.dset_gu, 0, NULL); - struct PC pc = {fmt, S, D, I, tg->rowWords, tg->gs}; // PC.I = input D, PC.O = moe_inter I + struct PC pc = {fmt, S, D, I, tg->rowWords, tg->gs, 0, 0.f, 0.f}; // PC.I = input D, PC.O = moe_inter I pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdPushConstants(G.cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G.cmd, (uint32_t)((I + 7) / 8), (uint32_t)S, 1); @@ -799,7 +799,7 @@ static int eg_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *up /* phase 1: fused gate+up+silu -> hidden (per expert, bound to its x/hidden slices) */ vkCmdBindPipeline(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); for (int c = 0; c < count; c++) { - struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs}; + struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs, 0, 0.f, 0.f}; pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdBindDescriptorSets(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gu, 0, 1, &G.eg_gu[c], 0, NULL); vkCmdPushConstants(G.eg_cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); @@ -809,7 +809,7 @@ static int eg_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *up /* phase 2: down projection hidden -> y */ vkCmdBindPipeline(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); for (int c = 0; c < count; c++) { - struct PC pc = {dfmt, rows[c], I, D, downs[c]->rowWords, downs[c]->gs}; + struct PC pc = {dfmt, rows[c], I, D, downs[c]->rowWords, downs[c]->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.eg_cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.eg_dn[c], 0, NULL); vkCmdPushConstants(G.eg_cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G.eg_cmd, (uint32_t)((D + 7) / 8), (uint32_t)rows[c], 1); @@ -1142,7 +1142,7 @@ static int eg2_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *u VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; vkCmdBindPipeline(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.pipe_gu); for (int c = 0; c < count; c++) { - struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs}; + struct PC pc = {fmt, rows[c], D, I, gates[c]->rowWords, gates[c]->gs, 0, 0.f, 0.f}; pc.act = vk_act; pc.alpha = vk_alpha; pc.limit = vk_limit; // gate activation (M3: swigluoai) vkCmdBindDescriptorSets(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.plyt_gu, 0, 1, &G2.gu[c], 0, NULL); vkCmdPushConstants(G2.cmd, G2.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); @@ -1151,7 +1151,7 @@ static int eg2_prepare_submit(ColiVkTensor *const *gates, ColiVkTensor *const *u vkCmdPipelineBarrier(G2.cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, NULL, 0, NULL); vkCmdBindPipeline(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.pipe); for (int c = 0; c < count; c++) { - struct PC pc = {dfmt, rows[c], I, D, downs[c]->rowWords, downs[c]->gs}; + struct PC pc = {dfmt, rows[c], I, D, downs[c]->rowWords, downs[c]->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G2.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G2.plyt, 0, 1, &G2.dn[c], 0, NULL); vkCmdPushConstants(G2.cmd, G2.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); vkCmdDispatch(G2.cmd, (uint32_t)((D + 7) / 8), (uint32_t)rows[c], 1); @@ -1388,7 +1388,7 @@ int coli_vk_gqa_attn_project(float *out, const float *q, ColiVkTensor **ot, cons VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, NULL, 0, NULL); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); - struct PC opc = {ofmt, S, H * hd, Dout, to->rowWords, to->gs}; + struct PC opc = {ofmt, S, H * hd, Dout, to->rowWords, to->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(opc), &opc); vkCmdDispatch(G.cmd, (uint32_t)((Dout + 7) / 8), (uint32_t)S, 1); VKCHECK(vkEndCommandBuffer(G.cmd), "endCmd"); @@ -1442,11 +1442,11 @@ int coli_vk_matmul_pair(ColiVkTensor **t1p, float *y1, const void *w1, const flo VkCommandBufferBeginInfo begin = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; VKCHECK(vkBeginCommandBuffer(G.cmd, &begin), "beginCmd"); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); - struct PC pc1 = {fmt, S, I, O1, t1->rowWords, t1->gs}; + struct PC pc1 = {fmt, S, I, O1, t1->rowWords, t1->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc1), &pc1); vkCmdDispatch(G.cmd, (uint32_t)((O1 + 7) / 8), (uint32_t)S, 1); - struct PC pc2 = {fmt, S, I, O2, t2->rowWords, t2->gs}; + struct PC pc2 = {fmt, S, I, O2, t2->rowWords, t2->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset_pair, 0, NULL); vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc2), &pc2); vkCmdDispatch(G.cmd, (uint32_t)((O2 + 7) / 8), (uint32_t)S, 1); @@ -1529,11 +1529,11 @@ int coli_vk_attn_qprep(int layer, VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); - struct PC pc1 = {fmt, S, I, Oqa, tqa->rowWords, tqa->gs}; + struct PC pc1 = {fmt, S, I, Oqa, tqa->rowWords, tqa->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc1), &pc1); vkCmdDispatch(G.cmd, (uint32_t)((Oqa + 7) / 8), (uint32_t)S, 1); - struct PC pc2 = {fmt, S, I, Okva, tkv->rowWords, tkv->gs}; + struct PC pc2 = {fmt, S, I, Okva, tkv->rowWords, tkv->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset_pair, 0, NULL); vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc2), &pc2); vkCmdDispatch(G.cmd, (uint32_t)((Okva + 7) / 8), (uint32_t)S, 1); @@ -1547,7 +1547,7 @@ int coli_vk_attn_qprep(int layer, vkCmdPipelineBarrier(G.cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, NULL, 0, NULL); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); - struct PC pc3 = {fmt, S, Oqa, Oqb, tqb->rowWords, tqb->gs}; + struct PC pc3 = {fmt, S, Oqa, Oqb, tqb->rowWords, tqb->gs, 0, 0.f, 0.f}; vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset_qp3, 0, NULL); vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc3), &pc3); vkCmdDispatch(G.cmd, (uint32_t)((Oqb + 7) / 8), (uint32_t)S, 1); @@ -1621,7 +1621,7 @@ int coli_vk_attention_absorb_project(ColiVkTensor **kvb, const void *w, const fl VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, NULL, 0, NULL); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); - struct PC opc = {ofmt, S, H * V, Dout, to->rowWords, to->gs}; + struct PC opc = {ofmt, S, H * V, Dout, to->rowWords, to->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(opc), &opc); vkCmdDispatch(G.cmd, (uint32_t)((Dout + 7) / 8), (uint32_t)S, 1); VKCHECK(vkEndCommandBuffer(G.cmd), "endCmd"); @@ -1831,7 +1831,7 @@ static double bench_batched(ColiVkTensor *t, const float *x, int fmt, int S, int vkBeginCommandBuffer(G.cmd, &begin); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt, 0, 1, &G.dset, 0, NULL); - struct PC pc = {fmt, S, I, O, t->rowWords, t->gs}; + struct PC pc = {fmt, S, I, O, t->rowWords, t->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; @@ -1887,7 +1887,7 @@ static double bench_gu_batched(ColiVkTensor *tg, const float *x, int fmt, int S, vkBeginCommandBuffer(G.cmd, &begin); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); vkCmdBindDescriptorSets(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.plyt_gu, 0, 1, &G.dset_gu, 0, NULL); - struct PC pc = {fmt, S, D, I, tg->rowWords, tg->gs}; + struct PC pc = {fmt, S, D, I, tg->rowWords, tg->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; @@ -1935,7 +1935,7 @@ static double bench_experts_fair(int fmt, int D, int I, int K, int Npass) { VkCommandBufferBeginInfo begin = {.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO}; vkBeginCommandBuffer(G.cmd, &begin); vkCmdBindPipeline(G.cmd, VK_PIPELINE_BIND_POINT_COMPUTE, G.pipe_gu); - struct PC pc = {fmt, 1, D, I, tg[0]->rowWords, tg[0]->gs}; + struct PC pc = {fmt, 1, D, I, tg[0]->rowWords, tg[0]->gs, 0, 0.f, 0.f}; vkCmdPushConstants(G.cmd, G.plyt_gu, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(pc), &pc); VkMemoryBarrier mb = {.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER, .srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT, .dstAccessMask = VK_ACCESS_SHADER_READ_BIT}; for (int pass = 0; pass < Npass; pass++) for (int c = 0; c < K; c++) { diff --git a/c/colibri.c b/c/colibri.c index 9dd42f5c0..35d66285e 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -7442,7 +7442,9 @@ static void kv_alloc(Model *m, int max_t){ k->Lc8=k->Rc8=NULL; k->Lsc=k->Rsc=NULL; } if(k->Ic){ for(int i=0;in_layers;i++) free(k->Ic[i]); free(k->Ic); k->Ic=NULL; } if(m->has_dsa || c->msa){ /* index-key cache: GLM DSA or M3 MSA */ - k->Ic=calloc(c->n_layers,sizeof(float*)); + /* (unsigned): n_layers is config-validated positive; the cast gives VRP a + * provable bound so -Walloc-size-larger-than stays quiet on inlined paths. */ + k->Ic=calloc((unsigned)c->n_layers,sizeof(float*)); for(int i=0;in_layers;i++) if(c->idx_type[i]) k->Ic[i]=falloc((int64_t)max_t*c->index_hd); } k->max_t=max_t; diff --git a/docs/minimax-m3.md b/docs/minimax-m3.md new file mode 100644 index 000000000..32ffdaefc --- /dev/null +++ b/docs/minimax-m3.md @@ -0,0 +1,75 @@ +# MiniMax-M3 (rides `c/colibri.c`) + +Support for [MiniMax-M3](https://huggingface.co/MiniMaxAI/MiniMax-M3) +(426B parameters, 60 layers, 128 experts / top-4 + 1 shared) inside the GLM +engine — **not** a sibling binary. `colibri.c` reads `config.json` at startup +and switches to `ARCH_M3`; the control plane says so once, in the +`minimax_m3` `FamilyDescriptor` (`c/family_registry.py`), which declares +`engine_artifact="colibri"` and `engine_group="colibri-core"` — the same +group as GLM, and the reason `coli` hands M3 the GLM binary's env contract, +CAP channel and byte-protocol REPL rather than treating it as a sister +engine. Everything downstream — serve, web, chat, tools, the Vulkan expert +tier — is the GLM plumbing. + +``` +python3 c/tools/convert_fp8_to_int4.py --arch m3 --repo MiniMaxAI/MiniMax-M3 \ + --outdir m3_i4 --ebits 4 --io-bits 8 # --arch auto also detects it +./c/coli chat --model m3_i4 +./c/coli serve --model m3_i4 # OpenAI API, tool calling included +``` + +## Architecture notes + +Relative to GLM-5.2 (MLA + DSA), M3 swaps in four pieces: + +- **GQA attention.** Plain q/k/v projections, 64 query heads over 4 KV heads, + explicit `head_dim` 128 (≠ hidden/heads). Per-head **Gemma RMSNorm** + (`x/rms·(1+w)`) on Q and K *before* a partial split-half NEOX RoPE over the + first `rotary_dim` dims. KV cache is the full K and V rows per token — no + latent compression. +- **MSA (Lightning Indexer).** A small scoring branch (4 index heads, + dim 128) max-pools per-key scores into 128-token blocks and keeps the + top-16 blocks per query (+1 forced local block). Layers 0–2 are dense/full; + 3–59 are sparse. `COLI_MSA=0` disables selection for A/B runs (full causal + attention everywhere). +- **swigluoai activation.** `gate = min(g, limit)`, `up = clamp(u, ±limit)`, + `out = (up+1)·gate·σ(alpha·gate)` with `alpha`/`limit` from the config — + one dispatch point, `act_glu()`, shared with GLM's silu path. +- **Sigmoid router.** Raw sigmoid weights renormalized over the top-4 chosen + on bias-corrected scores, times `routed_scaling_factor`, plus the shared + expert. + +## Context and MSA exactness + +Two honest caveats worth knowing before raising `CTX`: + +- The engine default is `CTX=4096` — a *conservative* default for a model + whose selling point is long context. Raise it explicitly + (`CTX=32768 ./c/coli serve …`); `coli plan` sizes the KV reservation from + the real GQA + indexer-cache formulas, via the descriptor's + `_minimax_geometry` adapter. +- MSA selection is **exact** for attention windows up to + `sparse_topk_blocks × sparse_block_size` = 16 × 128 = **2048 tokens** (the + indexer selects every causal block then). Beyond that it is the model's own + trained approximation — that is how MiniMax runs it, but teacher-forcing + comparisons against a full-attention reference will diverge past 2048 by + construction. + +## Tool calling + +`coli serve` renders and parses the official `chat_template.jinja` dialect +byte-exactly: `]<]minimax[>[`-prefixed ``/`` XML blocks, +nested values via ``/child tags, tool results as `` messages. +Streaming holds output at the first tool-call marker instead of leaking the +raw block. EOS is `[e~[`, resolved by the shared `tok_eos_resolve()` for +`run`, `chat`, `serve` and batched serve alike. + +## Validation + +`make -C c m3-tiny-check` builds a tiny random M3 checkpoint, converts it +with `--arch m3`, runs the numpy oracle (`c/tools/oracle_m3.py`) and requires +the engine to match every teacher-forced prefill and decode position +token-exactly (CI runs this on every PR — it is the only gate that notices +when a GLM-side refactor of the shared hot paths breaks M3). `IDOT=0` for the +compare: int8 activation quantization can legitimately flip a borderline +argmax. From 479a4fc34e7a52fe8d31e13939991964dbb8b98a Mon Sep 17 00:00:00 2001 From: Steve Markgraf Date: Thu, 27 Aug 2026 14:03:52 +0200 Subject: [PATCH 22/22] docs(m3): correct the attention_gqa header, flag the M3 globals for Segment Review feedback from #601. - attention_gqa's header claimed "MSA block selection is NOT implemented: full causal attention". It has been implemented since the Lightning Indexer landed: `sel` is built per (row, index head) and gates both the score loop and the value accumulation, so only the selected blocks are read. The comment now says what the body does, keeps the exactness note (selection is a no-op while the window fits in topk*block = 2048 tokens, because every causal block is selected then) and names the three ways out of the sparse path: dense layers, no indexer weights, COLI_MSA=0. - g_gemma_norm / g_act_swigluoai / g_swiglu_{alpha,limit} are process globals, which held while a process opened one model. #1227 and #1245 landed the Segment and Edge runtimes, so a GLM segment can now sit beside an M3 segment and the two would fight over these: last model_init wins. Latent -- no shipping binary registers those adapters -- but M3's segment adapter cannot land until they hang off the Model. Recorded at the declaration, where the next person to add an adapter will read it. - colibri_core() resolved family_by_id("glm") on every call to reach a string constant, and would have raised if that id were ever renamed. The group name moves into family_registry as COLIBRI_CORE_GROUP and both descriptors reference it, so the shared engine is single-source and the predicate no longer depends on another family's id. --- c/coli | 7 ++++--- c/colibri.c | 19 +++++++++++++++---- c/family_registry.py | 9 +++++++-- 3 files changed, 26 insertions(+), 9 deletions(-) diff --git a/c/coli b/c/coli index bf8fb9bd6..8382a1534 100755 --- a/c/coli +++ b/c/coli @@ -68,8 +68,9 @@ except ModuleNotFoundError: except ModuleNotFoundError: _version = "unknown" -from family_registry import (FamilyConfigError, UnknownFamilyError, all_families, - family_by_id, resolve_model, tuning_replay_prompt) +from family_registry import (COLIBRI_CORE_GROUP, FamilyConfigError, UnknownFamilyError, + all_families, family_by_id, resolve_model, + tuning_replay_prompt) _TUNE_ROTATION_PROMPT = ( "A service has high median throughput but poor tail latency. Explain how " @@ -323,7 +324,7 @@ def colibri_core(family): id == "glm" so a future third family on this binary inherits the same handling by declaring the group, not by being added to a tuple here.""" if isinstance(family, str): family = family_by_id(family) - return family.engine_group == family_by_id("glm").engine_group + return family.engine_group == COLIBRI_CORE_GROUP def engine_for(model): family = resolve_model(model).descriptor diff --git a/c/colibri.c b/c/colibri.c index 35d66285e..e49bde79a 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -1508,7 +1508,13 @@ static void qt_fill(QT *t, const float *w, int bits){ } /* ARCH_M3 numeric conventions, set once at model_init from the config. Both default - * to the GLM behavior so every existing path is bit-identical when they stay 0. */ + * to the GLM behavior so every existing path is bit-identical when they stay 0. + * These are PROCESS-wide, which was safe while a process held one model. The Segment + * and Edge runtimes (#1227/#1245) can keep several engines open at once, and a GLM + * segment beside an M3 segment would fight over all four: last model_init wins and the + * loser silently gets the other's norm and activation. Latent today -- no shipping + * binary registers those adapters -- but M3's segment adapter cannot land until these + * hang off the Model rather than the process. */ static int g_gemma_norm=0; /* rmsnorm scales by (1+w) instead of w */ static int g_act_swigluoai=0; /* glu = clamp+alpha-sigmoid, (up+1)*glu */ static float g_swiglu_alpha=1.702f, g_swiglu_limit=7.0f; @@ -4251,9 +4257,14 @@ static void kv_lc_rows_f32(Model *m, int layer, int64_t t0, int64_t n, float *ds * KV rows in the Lc/Rc aliases (K = Lc, V = Rc, both n_kv_heads*head_dim wide — set * up by load_cfg); score/softmax/value with H/n_kv_heads query heads per KV head; * o_proj at the end. Mirrors attention_rows' ragged contract (kvs/positions per-row - * KV target, else pos_base+s). MSA block selection is NOT implemented: full causal - * attention — EXACT for windows up to sparse_topk_blocks*block = 2048 tokens (the - * indexer selects every block then), an approximation beyond. */ + * KV target, else pos_base+s). On sparse layers (idx_type[], for M3 exactly the MoE + * layers) the MSA Lightning Indexer picks the top-k key blocks per query and `sel` + * gates both the score loop and the value accumulation, so only the selected blocks + * are attended; dense layers, checkpoints without the indexer weights, and COLI_MSA=0 + * take the full causal path. Selection is a no-op while the window fits in + * sparse_topk_blocks*block = 2048 tokens (every causal block is selected then, so the + * key set is full attention's, summation order aside); past that it is the model's own + * trained approximation, which is how MiniMax runs it. */ static void rope_half_neox(float *v, int rot, int pos, float theta){ int h2=rot/2; for(int j=0;j