token.nu

packages/swarm-mcp/src/token.nu — cluster identity + authenticity from one shared --token.

Every node in a cluster is launched with the same --token. The token does two things, both derived deterministically so no node has to coordinate:

1. ISOLATION — token_group_id hashes the token into the 32-byte relay

multicast group id. Clusters with different tokens join different relay groups, so their HELLO/job traffic is mutually invisible even when they share one relay. Without the token you cannot derive the group id, so a stranger cannot even see (let alone join) the cluster's gossip.

1. AUTHENTICITY — token_key derives a 32-byte HMAC key. Every compute

payload (a kernel chunk) and every result is tagged with HMAC-SHA256 (token_tag) and verified on receipt (token_untag, constant-time). A worker only acts on token-authentic jobs; a coordinator only accepts token-authentic results. Combined with group isolation this gives mutual authentication of all compute traffic over the dumb (opaque) relay.

The relay itself never needs the token — it forwards opaque bytes — so the secret never leaves the nodes that own it.

API

@ token_tag_len → i

HMAC tag length prepended to authenticated messages (truncated SHA-256 MAC).

@ token_group_id s token → ( Vec u )

32-byte relay multicast group id for this cluster (token isolation).

@ token_key s token → ( Vec u )

32-byte HMAC key authenticating compute payloads + results.

@ token_tag ( Vec u ) key ( Vec u ) payload → ( Vec u )

Prepend a 16-byte HMAC-SHA256 tag to payload. Returns a NEW vec; the caller owns it (and still owns payload).

@ token_untag ( Vec u ) key ( Vec u ) tagged → ?( Vec u )

Verify the tag and strip it. Some(payload copy) on a valid tag, None on a missing/short/forged tag. The compare is constant-time. Caller owns the returned payload on the Some arm.

---

census.nu

packages/swarm/src/census.nu — lightweight membership gossip for the swarm.

The compute layer (dist/job over dist/ring) needs every node to agree on the live worker set: a worker must know it owns a key, and the coordinator must know which pubkeys to route to. The full SWIM table (net/membership.nu) is the heavy, churn-hardened answer; this is the small one a compute cluster actually needs — a HELLO announcement that feeds the consistent-hash ring.

• a node joins by broadcasting HELLO(role, want=1) to the relay group;
• everyone who hears it adds the node to their ring (workers only) and, if

want, replies with their own HELLO so the newcomer learns them too;

• the ring converges, and "join the cluster" is literally "run the binary".

Roles: a WORKER owns keys and executes handlers, so it joins everyone's ring; a CLIENT (the coordinator) only submits, so it is never added to the ring — otherwise it would own keys it has no handler for and silently drop results.

The codec and the membership set are pure; only the pump touches transport.

API

@ census_hello_t → i

@ role_client → i

@ role_worker → i

@ cap_gpu → i

Capability bits a worker advertises in HELLO. cap_gpu: the worker runs wasm chunks with GPU host imports enabled (nwasm --allow-gpu on real hardware).

@ hello_build i id i role i want ( Vec u ) pubkey i caps → ( Vec u )

HELLO wire: [3][id:8][role:1][want:1][pklen:2][pubkey…][caps:1] The caps byte TRAILS the pubkey so a pre-caps decoder (fixed prefix + pklen-delimited pubkey) reads the same fields and ignores the tail; a caps decoder treats a missing tail as caps=0. Mixed-version clusters stay sound.

: Hello { i id i role i want ( Vec u ) pubkey i caps }

@ hello_free Hello h → v

@ hello_decode ( Vec u ) buf → Hello

: Member { ( Vec u ) pubkey i id i caps i last_ms }

last_ms is when this member's HELLO was last heard. Workers re-announce on a ~2 s heartbeat, so a member silent far longer than that is gone: without this the ring kept a dead worker forever and every later submit paid a full round of chunk re-dispatches routing at a node that is never coming back. Eviction is self-healing — a worker that comes back re-announces and rejoins.

: Roster { ( Vec s ) members } // *Member

@ roster_new → *Roster

@ roster_free * Roster r → v

@ roster_has * Roster r ( Vec u ) pubkey → b

@ roster_count * Roster r → i

@ roster_count_caps * Roster r i mask → i

How many members advertise every capability bit in mask.

@ roster_add * Roster r * Ring ring ( Vec u ) pubkey i id i vnodes i caps i now → b

Fold a worker into the roster + ring, once. Returns T if newly added. now is the caller's clock (ms); the member's liveness stamp starts there.

@ roster_touch * Roster r ( Vec u ) pubkey i now → v

Refresh a member's liveness stamp (a re-heard HELLO). Unknown pubkey: no-op.

@ roster_expire * Roster r i now i ttl_ms ( Vec u ) exempt → ( Vec ( Vec u ) )

Drop every member silent for longer than ttl_ms and return their pubkeys (owned by the caller, which also removes them from its rings). The roster entry — not the ring — is the source of truth for who is live, so callers must apply the returned removals to every ring they maintain.

exempt is never evicted (pass a node's own pubkey, empty for none): a worker hears no HELLO of its own, so without the exemption it would time itself out of its own ring and stop owning — and therefore stop executing — every key it holds.

@ roster_is_live * Roster r ( Vec u ) pubkey → b

True when pubkey is a live roster member (the coordinator's liveness test for the worker a chunk was routed to).

: MemberView { i id i caps i last_ms }

Read-only view of member k: its node id, capability bits and last-heard stamp. Out of range → id 0. Used by the status tool, so an operator (or the model) can see the cluster the coordinator believes it has.

@ roster_view * Roster r i k → MemberView

---

work.nu

packages/swarm-mcp/src/work.nu — the distributed map-reduce the cluster runs.

A task is: evaluate an expression kernel (expr.nu) over an integer range [lo, hi) and fold the results with a reduce op. The coordinator shards the range; dist/ring routes each chunk to its owning worker; the worker parses the kernel once and folds it over its sub-range; the coordinator combines the partial folds. Every reduce op is associative, so sharding is exact.

reduce op : 0 sum · 1 product · 2 min · 3 max · 4 count (of truthy map) dtype : 0 int (i64) · 1 float (f64) chunk : [op:1][dtype:1][lo:8 BE][hi:8 BE][expr bytes…] result : [value:8 BE] — i64, or an f64 bit pattern when dtype=float

Float tasks fold in f64 over the same integer range (x is the index cast to double) and ship the partial as its f64 bit pattern through the same 8-byte result codec; the coordinator reinterprets it (see main.nu tids_combine).

API

@ red_sum → i

@ red_product → i

@ red_min → i

@ red_max → i

@ red_count → i

@ red_id i op → i

The identity (empty-range) value for a reduce op.

@ red_fold i op i acc i v → i

Fold one mapped value into the accumulator (per element).

@ red_combine i op i a i b → i

Combine two partial folds (across chunks). count combines by summing the per-chunk counts; every other op combines like its element fold.

@ red_id_f i op → f

@ red_fold_f i op f acc f v → f

@ red_combine_f i op f a f b → f

@ chunk_payload i op i dtype i lo i hi ( Vec u ) expr → ( Vec u )

@ result_encode i value → ( Vec u )

@ result_decode ( Vec u ) p → i

@ kernel_handler ( Vec u ) key → ( @ ( Vec u ) ( Vec u ) )

@ kernel_handler_ka ( Vec u ) key ( @ v ) ka → ( @ ( Vec u ) ( Vec u ) )

: Chunk { i lo i hi }

@ shard i lo i hi i n → ( Vec s )

@ shard_free ( Vec s ) chunks → v

@ chunk_key i idx → ( Vec u )

A ring key for chunk index i: distinct chunks hash to distinct ring points.

@ reduce_op_known s name → b

True for a recognised reduce-op name. The MCP layer rejects anything else instead of quietly folding with the fallback: an unknown op used to return a plausible-but-wrong number ("avg" reduced as sum), which is the one failure an agent cannot detect.

@ reduce_op_of s name i fallback → i

True for a recognised reduce-op name; sets *op. Keeps the MCP layer thin.

@ reduce_op_name i op → s

---

expr.nu

packages/swarm-mcp/src/expr.nu — the phase-1 "kernel": a small, regular integer-expression language the cluster evaluates per element.

A workload is a map-reduce: the coordinator ships an expression in one variable x plus a range and a reduce op; each worker parses the expression once, evaluates it for every x in its sub-range, and folds the results. The language is deliberately small and regular so a language model can write it without docs — and it is the natural precursor to phase 2, where the kernel is arbitrary NURL compiled to wasm instead of interpreted here.

expr := ternary ternary := logic ( '?' expr ':' ternary )? logic := compare ( ('&'|'|') compare ) compare := addsub ( ('<'|'<='|'>'|'>='|'=='|'!=') addsub )? addsub := muldiv ( ('+'|'-') muldiv ) muldiv := unary ( (''|'/'|'%') unary ) unary := '-' unary | primary primary := NUM | 'x' | '(' expr ')' | ('min'|'max') '(' expr ',' expr ')' | 'abs' '(' expr ')' NUM := DIGIT+ ( '.' DIGIT+ )?

The same grammar evaluates in one of two numeric domains, chosen per task by the caller (see work.nu dtype): • int — all arithmetic is i64 (truncated division; div/mod by zero → 0). • float — all arithmetic is f64, x is the integer index cast to double, div/mod by zero → 0.0. A literal with a '.' (e.g. 0.5) is a float literal in either mode (truncated toward zero in int mode). Comparisons and '&' '|' yield 1/0 (1.0/0.0 in float); any non-zero is "true".

Tokens are kept in two parallel int vectors; the parser builds a flat stride-4 node arena (tag,a,b,c) and returns the root index — the resp.nu arena pattern, so nesting needs no per-node allocation. Float literals carry the f64 bit pattern (via floatbits) in the value slot, so the all-int arena holds them losslessly; the float evaluator reinterprets them back.

API

@ expr_tokenize ( Vec u ) src ( Vec i ) tk ( Vec i ) tv → b

: EParser

: EParser {
    ( Vec i ) tk
    ( Vec i ) tv
    i pos
    ( Vec i ) arena  // stride-4: tag,a,b,c
    b ok
}

@ expr_parse ( Vec u ) src * EParser p → i

Parse src → (arena, root). On any error, ok=0; the caller checks it. The arena is returned via the EParser; the root index via the return value.

@ eparser_free * EParser p → v

Frees the parser's arena/token vectors AND the struct itself.

@ expr_eval * EParser p i node i x → i

@ expr_eval_f * EParser p i node f x → f

---

buildwasm.nu

packages/swarm-mcp/src/buildwasm.nu — compile NURL source to a wasm module so the MCP server itself can accept a kernel as source (compute_submit_kernel / compute_submit_cuda) instead of requiring the caller to pre-compile.

LOCAL-FIRST: the wasmbuilder package (deps/wasmbuilder) compiles the kernel in-process — nurlc → IR rewrite → the toolchain's bundled zig cc — no network, no build service. Only when the local toolchain can't do it (no nurlc/zig on this box) does it fall back to POSTing {source, filename} to <NURL_BUILD_API>/build_wasm. A direct nurlapi answers with raw wasm bytes; the public playground proxy answers with JSON carrying wasm_base64 (and nurlc_errors on failure). Both are handled.

$NURL_BUILD_API fallback build service base URL (default https://play.nurl-lang.org)

API

@ build_api_url → String

@ wrap_kernel s source i op i dtype i kkind → String

Wrap a bare kernel into a complete wasm-ready program for reduce op op. dtype 0 (int): kernel is @ kernel i x → i; the module folds in i64 and prints the partial as a decimal integer. dtype 1 (float): kernel is @ kernel i x → f (x is the integer index, returns a double); the module folds in f64 and prints the partial's f64 BIT PATTERN as a decimal integer — so it rides the same stdout→int wire, and the coordinator reinterprets it (work.nu tids_combine float path).

kkind 0 (element): the generated main folds kernel(x) over [lo, hi). kkind 1 (chunk): the kernel is @ kernel i lo i hi → i (or → f) and the main calls it ONCE with the whole sub-range — the kernel owns the loop. This is the right granularity for kernels with per-invocation setup cost (open a CUDA context, JIT a device kernel, allocate buffers): one setup per CHUNK instead of one per element. The reduce op still combines the chunk partials.

@ compile_to_wasm s source → !( Vec u ) String

Compile NURL source to a wasm module. Ok = module bytes; Err = a human-readable transport/compile error (suitable to hand back to the model).

Local wasmbuilder first. A LOCAL "nurlc failed" is a genuine source error and is returned as-is (the build service would only repeat it); any other local error means the environment can't build wasm (no toolchain, no zig, download forbidden) → fall back to the build API.

---

cudakernel.nu

packages/swarm-mcp/src/cudakernel.nu — generate a GPU chunk kernel from a bare CUDA-C map function.

The model hands over ONLY the math; cuda_wrap wraps it into a complete, self-contained NURL program in one of three OUTPUT MODES:

scalar (gpu_mode_scalar) — __device__ double f(long long x): grid-stride map over [lo, hi), per-block shared-memory reduction with the task's reduce op, host fold of the block partials, partial's f64 bit pattern printed on stdout (the classic reduce wire). sample (gpu_mode_sample) — same f, but every value comes BACK: out[x-lo] = f(x); the module writes the hi-lo doubles raw (LE) to an output file named on argv — curves, fields, images. hist (gpu_mode_hist) — __device__ long long bin(long long x) plus an optional __device__ double val(long long x) (default 1.0): atomicAdd(&out[bin(x)], val(x)) over K bins (K on argv — the same module serves any K), K doubles written to the output file. The double atomicAdd is a portable CAS loop, so the PTX stays valid for NVRTC's default (pre-sm_60) target.

RUNTIME PARAMS make the kernels dynamic: with has_params the device functions take (long long x, const double* p) and the module parses any argv tail of f64-bit-pattern decimals into a device buffer. Parameter VALUES never touch the generated source, so the module's content hash — and with it the worker-side cache and the coordinator's build cache — survives across a whole parameter scan: no rebuild, no re-upload.

The program's cuda/nvrtc FFI declarations mirror packages/gpu/src/cuda.nu's ABI exactly. Compiled with --ffi-host-imports (the wasm build API always does), every cu/nvrtc symbol becomes an env import that the pure-NURL nwasm's GPU bridge resolves against the worker's real libcuda/libnvrtc under --allow-gpu. The same source also compiles native (nurl.sh auto-links cuda/nvrtc), so a kernel can be smoke-tested off-cluster.

Argv contracts (must match wasmkernel.nu __wasm_run_gpu): scalar: main lo hi [p…] sample: main lo hi outpath [p…] hist: main lo hi outpath K [p…]

Any CUDA failure exits non-zero WITHOUT producing output — the worker reports the chunk as failed (ok=0) instead of folding silent zeros.

API

@ cuda_src_ok s src → b

── validation ──────────────────────────────────────────────────── The user source is spliced into a generated NURL backtick literal; a backtick would terminate it (NURL strings cannot contain one, escaped or not), so reject it outright. Everything else is escaped below.

@ cuda_wrap s user i op i mode i has_params i has_data → String

── the wrapper ─────────────────────────────────────────────────── CUDA map function + reduce op + output mode (+ params flag) → complete NURL chunk-kernel program. The caller validates cuda_src_ok first. has_data: the chunk carries a dataset slice — argv gains an in-file path right after hi, and the device functions receive v = data[x] as their second argument.

@ cuda_wrap_update s user i S i A → String

── the update kernel (compute_iterate's general step rule) ─────────── The iteration engine's per-round update — new_state[j] = update(j, …) — is nothing but a sample over [0, S): the model's update() runs once per state component with the whole state, the reduced accumulator, N and the runtime params in scope. It rides the ordinary sample mode; we only wrap the model's update() in the sample interface (f(x, p)) and expose named accessors into the packed param buffer p = [state (S) | acc (A) | N | params…].

device double update(long long j, const double* p) swarm_state(i) — the current parameter vector, p[0..S) swarm_acc(i) — the reduced accumulator from grad(), p[S..S+A) swarm_N — the element count over the range/dataset swarm_param(i) — the runtime params (learning rate, etc.) swarm_dim/swarm_adim — S and A, for looping (norms, per-group reduce)

---

main.nu

packages/swarm-mcp/src/main.nu — swarm-mcp: an MCP-controlled distributed compute engine. A language model sets a workload over MCP — an expression kernel in x plus a range and a reduce op — and the cluster evaluates it distributed; the model reads running tasks and finished results back.

swarm-mcp relay 0.0.0.0 47700 [--v] # the meeting point (one per cluster) swarm-mcp worker <host> <port> # join as a compute node swarm-mcp mcp <host> <port> # MCP server (stdio) → drives the cluster swarm-mcp submit <host> <port> <reduce> <lo> <hi> <expr> # manual CLI submit

The cluster layer (membership → ring → job dispatch) is the swarm package's; here every worker registers ONE generic kernel handler (work.nu) that interprets the submitted expression (expr.nu), so any normal-operation workload runs without recompiling a worker. Phase 2 swaps the interpreter for a NURL→wasm kernel; the protocol and MCP surface stay the same.

API

@ swarm_vnodes → i

@ kind_kernel → i

@ pk_from_id i id → ( Vec u )

A 32-byte opaque routing pubkey derived deterministically from a node id.

: Swarm

: Swarm {
    s transport  // *Transport
    s ring  // *Ring
    s gpu_ring  // *Ring — the GPU capability domain (cap_gpu workers only)
    s roster  // *Roster
    s job  // *JobNode
    ( Vec u ) self_pk
    i self_id
    i role
    i self_caps  // capability bits this node advertises (cap_gpu for --gpu workers)
    ( Vec u ) group  // token-derived 32-byte relay multicast group (cluster isolation)
    ( Vec u ) key  // token-derived HMAC key (compute authenticity)
    i epoch  // bumps on every ring-membership change (block-seed invalidation)
}

@ swarm_new RelayClient rc i id i role i caps s token → *Swarm

@ swarm_free * Swarm sw → v

@ swarm_join_group * Swarm sw → v

@ swarm_announce * Swarm sw i want → v

@ swarm_announce_ok * Swarm sw i want → b

Announce presence; returns whether the broadcast reached the relay. A failed send is the reconnect loop's signal that the relay is gone.

@ swarm_on_hello * Swarm sw Hello h → v

@ swarm_expire * Swarm sw → i

Drop workers that stopped announcing, from the roster and from both rings. Returns how many were evicted; a non-zero result bumps the epoch, because a changed ring re-homes chunk keys and invalidates recorded block seeds.

@ swarm_pump * Swarm sw i max → v

@ swarm_discover * Swarm sw i rounds → v

@ parse_relay_list String cs s dhost i dport → ( Vec String )

Dial the relay, retrying briefly so a co-located relay that is still binding (or a relay that restarts) doesn't lose its local roles to a startup race. ── relay endpoint list (failover) ─────────────────────────────── --connect accepts a comma-separated list "h1:p1,h2:p2,…"; any one being reachable bootstraps the node, and if the connected relay dies the node rotates to the next. A single endpoint is just a one-element list. The list is kept as raw "host:port" String segments (default filled in when a segment omits either); parse_hostport resolves one at dial time.

@ relay_list_free ( Vec String ) lst → v

@ relay_dial_retry s host i port i tries → !RelayClient NetErr

@ node_relay s host i port i vflag → v

RELAY role: the dumb rendezvous point. Owns the fiber reactor on this thread.

@ node_worker ( Vec String ) relays s dhost i dport s token i vflag i gpu → v

WORKER role: join the cluster, register the kernel + wasm handlers (every worker can run wasm modules), and drain cluster traffic forever. Each worker thread takes a fresh random identity, so --workers N spins up N independent ring members in one process. gpu≠0: advertise cap_gpu and register the kind_wasm_gpu handler — wasm chunks of that kind run with --allow-gpu. A worker keeps its identity across reconnects (same ring member) and, when the relay it is on dies, rotates to the next relay in the list and re-forms — no single relay is a point of failure. Death is detected by a periodic heartbeat announce whose SEND fails when the relay is gone; the on-disk block cache survives the switch, so re-seeding is idempotent.

@ cluster_submit * Swarm sw i op i dtype i lo i hi ( Vec u ) expr i nchunks → ( Vec i )

@ cluster_submit_wasm * Swarm sw i lo i hi ( Vec u ) wasm i nchunks i kind → ( Vec i )

Shard a wasm-kernel task: ship the compiled module + each sub-range to its ring owner under kind (kind_wasm, or kind_wasm_gpu — the GPU capability domain). The module bytes ride every chunk (workers cache by content hash, so it is written once per worker).

: ChunkJob

: ChunkJob {
    i kind  // job kind (kind_wasm_gpu)
    ( Vec u ) payload  // tagged, immutable across retries
    i idx  // chunk index (base of the ring key)
    i salt  // key salt — bumped to re-route away from a failed owner
    i tid  // current attempt's job task id
    ( Vec u ) owner  // the owner this attempt routed to (empty if unknown)
    i attempts  // dispatches made (1 = initial)
    i submit_ms  // wall-clock ms of the current dispatch (deadline base)
    i state  // 0 pending · 1 ok · 2 exhausted (failed after max attempts)
}

@ chunk_key_salted i idx i salt → ( Vec u )

Ring key for chunk idx under retry salt. Both idx and salt are folded into a bit-MIXED 64-bit value (multiply by large odd constants — a Fibonacci / splitmix pair) so every (idx, salt) pair lands at a well-separated ring point. A plain [idx][salt] byte layout does NOT work: the ring hash (FNV-1a) has poor avalanche on trailing bytes, so consecutive salts would map to the same arc and a re-dispatch could never escape the failed owner.

@ cj_tids ( Vec s ) jobs → ( Vec i )

@ chunkjobs_free ( Vec s ) jobs → v

@ cluster_dispatch_gpu_ft * Swarm sw i mode i lo i hi i kbins ( Vec i ) params ( Vec u ) wasm i nchunks → ( Vec s )

Dispatch a non-dataset GPU task WITH a retry plan: same wire as cluster_submit_wasm_gpu, but each chunk keeps its tagged payload and owner so task_refresh can re-dispatch it. Returns the *ChunkJob vector (the caller derives tids with cj_tids and stores the plan on the Task).

@ cluster_dispatch_kernel_ft * Swarm sw i op i dtype i lo i hi ( Vec u ) expr i nchunks → ( Vec s )

Dispatch an EXPRESSION task with a retry plan. Same wire as cluster_submit; the plan is what lets task_refresh notice a worker that died mid-chunk. Without it an expression task whose owner disappeared stayed running forever, which an agent can only poll into infinity.

@ cluster_dispatch_wasm_ft * Swarm sw i lo i hi ( Vec u ) wasm i nchunks i kind → ( Vec s )

Dispatch a wasm-module task (CPU kind_wasm or GPU kind_wasm_gpu) with a retry plan — the module bytes ride each chunk exactly as cluster_submit_wasm.

@ cluster_submit_wasm_gpu * Swarm sw i mode i lo i hi i kbins ( Vec i ) params ( Vec u ) data ( Vec u ) wasm i nchunks → ( Vec i )

Shard a GPU task (payload v2/v3): mode, K, runtime params and the module ride every chunk under kind_wasm_gpu — routed on the GPU capability ring. data is the WHOLE dataset's raw LE f64 bytes (empty when the task has no dataset); each chunk ships exactly its own slice data[clo·8, chi·8) — the split travels with its task, so a worker needs no separate fetch.

@ tids_ready * Swarm sw ( Vec i ) tids → b

All chunk results present? (cluster job results are recorded by task-id.)

: Combined { i value i nfail }

Combine all chunk results with the reduce op (assumes ready). For a float task each partial is an f64 bit pattern: decode→f64, combine in f64, and return the combined f64 re-encoded as its bit pattern (the Task stores it that way; task_to_json reinterprets it for the model).

Two result frames ride the wire: the expr kernel's legacy [partial:8], and the wasm kinds' [ok:1][partial:8]. A chunk whose ok=0 (module trap, missing runtime, GPU failure) is COUNTED, not folded — nfail>0 marks the task as failed instead of silently reducing zeros into the answer.

@ tids_combine * Swarm sw i dtype i op ( Vec i ) tids → Combined

: CombinedV { ( Vec u ) bytes i nfail }

Combine VECTOR results (sample / hist). Frame: [ok:1][count:4 BE][f64 LE × count]. Sample chunks CONCATENATE in shard order (tids order = chunk order); hist chunks add ELEMENTWISE into K bins. A failed or malformed chunk counts toward nfail and contributes nothing.

@ tids_combine_vec * Swarm sw i mode i kbins ( Vec i ) tids → CombinedV

@ nchunks_for i nworkers → i

@ run_submit s host i port i op i lo i hi s expr s token → i

@ nchunks_wasm i nworkers → i

Chunk count for a wasm task: fewer chunks than the expression path (each chunk ships the module and spawns a runtime process), but enough to spread across the ring.

@ run_runwasm s host i port i op i lo i hi s wasmpath s token → i

: Task

: Task {
    i id
    ( Vec u ) expr  // kernel bytes
    i dtype  // 0 int · 1 float (result holds an f64 bit pattern)
    i lo
    i hi
    i op
    i nchunks
    ( Vec i ) tids
    i done
    i result
    i failed  // chunks that reported ok=0 (task status "error" when > 0)
    i mode  // gpu_mode_scalar/sample/hist (scalar for every non-GPU task)
    i kbins  // histogram bin count (hist mode)
    ( Vec u ) vres  // vector result bytes (raw LE f64s; sample/hist modes)
    String out_file  // when set, the finished vector result is written here
    i dsid  // dataset id the task maps over (0 = none)
    i seeded  // dataset blocks seeded by THIS submit (0 = all were cached)
    ( Vec s ) chunkjobs  // *ChunkJob — per-chunk retry plan (empty = no auto-retry)
    i retries  // total chunk re-dispatches performed (fault tolerance)
    String errmsg  // why the first failed chunk failed ("" = none / not failed)
}

: McpState

: McpState {
    s swarm  // *Swarm coordinator
    ( Vec s ) tasks  // *Task
    i next_id
    ( Vec s ) wcache  // *WasmCached — compiled kernel modules by source hash
    ( Vec s ) datasets  // *Dataset — uploaded data the CUDA tools map over
    i next_ds
    ( Vec String ) seeded  // "blockhex|chunk|epoch" — blocks CONFIRMED cached at their owner
    McpTaskStore mcptasks  // io.modelcontextprotocol/tasks handles, keyed by task id
    s last_task  // *Task the tool call currently in flight registered (0 = none)
}

@ task_new i id ( Vec u ) expr i dtype i lo i hi i op i nchunks ( Vec i ) tids → *Task

Constructor. The params share the field names (lo, hi, …); = . t lo lo is a field store — the field-store/local-shadow miscompile this used to work around was fixed in the compiler (NURL v0.10.4).

: ~ i g_mcp 0

@ task_register * Task t → v

Register a freshly-submitted task and remember it as the one the in-flight tool call produced. Every submit path goes through here so the MCP tasks layer (below) can link its handle to the swarm task without each tool having to thread the pointer back out.

last_task is a single-slot handoff, which is sound because the MCP endpoint is a SERIAL accept loop (server_run): exactly one tool call is ever in flight. It is cleared before dispatch and read immediately after, so a tool that submits nothing leaves it at 0.

: Dataset { i id String name ( Vec u ) bytes ( Vec ( Vec u ) ) blocks String path i nbytes i dtype }

A dataset is either RAM-held (base64 upload → bytes) or FILE-BACKED (a path on the MCP host → path, bytes empty). Either way it is described by its content-address manifest (blocks) and total byte count (nbytes); the coordinator never needs the whole thing in RAM for a file-backed set — it streams block bytes from the file only when seeding a worker that lacks them.

@ ds_esz_of * Dataset d → i

dtype is the storage type code (== the kernel's has_data value): 1 f64 · 2 f32 · 3 i32 · 4 i64. Elements are stored in their native width; the count and all byte offsets scale by that width. The GPU promotes each to a double.

@ ds_count_of * Dataset d → i

@ ds_is_file * Dataset d → b

@ ds_dtype_id i dsid → i

The dtype code for a dataset id (1 f64 default if unknown) — the has_data value the CUDA generator and the chunk payload carry for this dataset.

@ ds_read_f ( Vec u ) bytes i off i dtype → f

Read one element at byte offset off as a double, honouring the dtype.

@ ds_block_bytes * Dataset d i b → ( Vec u )

Block b's raw bytes: sliced from RAM, or read from the source file at the block's fixed grid offset (file-backed). One block (≤ 1 MiB) at a time — the coordinator's memory never scales with the dataset size.

@ ds_file_manifest_stats s path i nbytes i dtype Json o → ( Vec ( Vec u ) )

Stream a file into its block manifest AND min/max/mean in ONE pass, holding at most one block (≤ 1 MiB) at a time. This is what lets a dataset be far larger than coordinator RAM: only the manifest (32 bytes/block) is retained.

@ ds_find i id → s

@ cluster_submit_wasm_gpu_ds * Swarm sw i mode i rlo i rhi i kbins ( Vec i ) params * Dataset d ( Vec u ) wasm i nchunks_want ( Vec String ) seeded * u nseed_cell → ( Vec i )

one-pass min/max/mean of raw LE f64 bytes, attached to a JSON object Dataset-backed GPU submit over CONTENT-ADDRESSED blocks (payload v4). Chunking is block-aligned, so every 1 MiB block belongs to exactly one chunk; blocks not yet confirmed at their owner are seeded first as kind_blob tasks carrying the SAME ring key as the compute chunk (same ring + same key → same worker), and per-connection ordering guarantees the worker handles a seed before the compute that references it. A seed is recorded in seeded (hash|chunk|epoch) only after an OK result, so a lost seed re-seeds on the next submit instead of wedging. Dataset-backed GPU submit over CONTENT-ADDRESSED blocks (payload v4). Chunking is block-aligned, so every block belongs to exactly one chunk. Blocks not yet confirmed at their owner are seeded FIRST, strictly one at a time — submit a block, pump until its OK result, then the next. A single large frame is ever in flight to a given worker, which keeps the relay's forwarding stream from interleaving multi-frame bursts. Only after every referenced block is confirmed cached are the (small) compute chunks submitted; each references its blocks by hash and the worker assembles the slice from its cache, failing visibly on any missing block.

: WasmCached { String hash ( Vec u ) wasm }

@ mcp_swarm → *Swarm

@ mcp_pump i rounds → v

@ task_refresh * Task t → v

@ task_find i id → s

@ task_to_json * Task t → Json

Build the JSON-as-text result body for one task (LLM-readable + parseable).

@ tool_result_json Json o → Json

@ tool_submit Json args → Json

@ tool_run_wasm Json args → Json

@ tool_submit_kernel Json args → Json

compute_submit_kernel: the server compiles NURL source → wasm itself (via the build API), then runs it — the model hands over a kernel as plain NURL.

: IterRun

: IterRun {
    i id
    i status  // 0 running · 1 done · 2 error
    b have_update
    ( Vec u ) wasm
    ( Vec u ) uwasm
    i S
    i A
    i N
    i rlo
    i rhi
    i dsid
    f lr
    f eps
    i rounds
    i rnd
    i ran
    i failed
    i total_seeded
    f last_delta
    b converged
    ( Vec i ) xparams
    ( Vec f ) state
    ( Vec f ) grad
    ( Vec f ) prev
}

── async iterate runs ──────────────────────────────────────────────── compute_iterate's loop can also run as a RUN OBJECT advanced in bounded slices by compute_iterate_status polls — the same drain-as-you-poll contract the submit tools follow, so a long training is never at the mercy of one HTTP call's timeout. The run keeps everything a round needs; the compiled modules are freed the moment the run finishes.

: IterRuns

: IterRuns {
    ( Vec s ) v
}

: ~ i g_iter_runs 0

: ~ i g_iter_next 1

@ tool_iterate Json args → Json

@ tool_iterate_status Json args → Json

Advance an async iterate run by a bounded time slice and report it.

@ tool_shuffle Json args → Json

@ tool_submit_cuda Json args → Json

@ tool_sample_cuda Json args → Json

compute_sample_cuda: every f(x) comes back — the cluster gathers the range into one array (curves, fields, images).

@ tool_hist_cuda Json args → Json

compute_histogram_cuda: bin(x) picks a bucket, val(x) the weight (default 1.0); K bin sums come back — a whole distribution in one pass.

@ tool_upload_data Json args → Json

@ tool_list_data Json args → Json

@ tool_status Json args → Json

── tool: swarm_status ─────────────────────────────────────────── What the coordinator believes the cluster is. Without this, "no workers found" or a task that keeps retrying gives a model nothing to reason about: it cannot see whether a worker ever joined, whether it is GPU-capable, or whether the node it is waiting on has gone silent.

@ tool_list Json args → Json

@ tool_result Json args → Json

@ mcp_task_store → McpTaskStore

@ ms_prop Json props s name s ty s desc → v

@ ms_schema_submit → Json

@ ms_schema_result → Json

@ ms_schema_empty → Json

@ ms_schema_run_wasm → Json

@ ms_schema_submit_kernel → Json

@ ms_schema_submit_cuda → Json

@ ms_schema_shuffle → Json

@ ms_schema_iterate → Json

@ ms_schema_iterate_status → Json

@ ms_schema_sample_cuda → Json

@ ms_schema_hist_cuda → Json

@ ms_schema_upload_data → Json

@ tool_help Json args → Json

@ ms_schema_help → Json

@ build_tools_list → ( Vec Json )

@ dispatch_tool s name Json args → Json

@ sm_version → s

Single source of truth for the server version — the MCP handshake, server/discover, and the --version banner all read this (the handshake had drifted to a stale hand-written 0.20.0).

@ handle_initialize Json id ? Json params → Json

@ handle_discover Json id → Json

server/discover — 2026-07-28 servers MUST implement this; also the stdio/HTTP backward-compat probe a dual-era client tries first.

@ finish_reply Json env b modern → ?Json

Decorate an outgoing response for a MODERN (per-request _meta) request: 2026-07-28 servers SHOULD identify themselves in each result's _meta. Legacy requests pass through untouched.

@ handle_ping Json id → Json

@ handle_tools_list Json id → Json

@ handle_tools_call Json id Json params b want_task → Json

want_task is the per-request tasks capability the client declared. It only ENABLES augmentation — the decision stays the server's, per tool (__mcp_task_eligible) and per call (a tool that registered no swarm task returns its result directly).

@ handle_unknown Json id s method → Json

@ dispatch Json req → ?Json

: ~ i g_mcp_relays 0 // *( Vec String ) as an address

The coordinator's relay endpoints and its current index — the submit path rebuilds the swarm against the next relay when the current one dies (mcp_ensure_relay below). Held as globals because the MCP handlers are top-level functions over module state, like the swarm itself.

: ~ i g_mcp_from 0 // next relay index to try

: ~ s g_mcp_token ``

: ~ i g_mcp_reconnects 0

@ mcp_reconnect → b

Build a fresh coordinator swarm on the next reachable relay and swap it into the McpState, preserving tasks / datasets / caches. Returns F when no relay in the list is reachable. Called from the submit path when the current relay looks dead, so a relay failure does not take the API down.

@ mcp_ensure_relay → b

Before a submit: if a heartbeat to the current relay fails, the relay is gone — reconnect to the next. Returns T once a live relay is in place.

@ node_mcp ( Vec String ) relays s rhost i rport s mcp_host i mcp_port s cert s key s token → v

@ arg_int i idx → i

@ arg_eq i idx s lit → b

@ has_flag s name → b

True if name appears anywhere in argv (a boolean role/verbose flag).

@ flag_val s name s deflt → String

The token following --name in argv, or a copy of deflt. Owned String.

@ flag_int s name i deflt → i

: HostPort { String host i port }

@ parse_hostport String hp s dhost i dport → HostPort

Parse "host:port", "host", ":port", or "" with host/port defaults.

@ relay_dial_list ( Vec String ) lst s dhost i dport i from * u idx_cell → !RelayClient NetErr

Dial the relays in lst starting at from (wrapping once around the whole list), a few quick retries each. Writes the connected index to idx_cell. Returns the live client or F when every endpoint is down.

@ ensure_cert s cert_path s key_path → b

Ensure a TLS cert+key exist at the given paths, auto-minting a self-signed EC P-256 pair in pure NURL if absent (std/x509_gen.nu — no openssl, no subprocess), so --mcp works out of the box; pass --tls-cert/--tls-key for a real cert. CN + SAN are localhost (matching the previous openssl invocation minus the IP SAN — MCP clients pin or use --insecure against a self-signed cert either way). Returns T once both files exist.

@ cli_submit → i

@ cli_runwasm → i

@ version → v

-v is already taken by --verbose, so the version flag is long-form only.

@ usage → v

@ run_node → i

@ main → i

---

blob.nu

packages/swarm-mcp/src/blob.nu — content-addressed dataset blocks.

The v3 payload shipped each chunk's dataset slice INSIDE every task, so N submits (or N iteration rounds) over the same data paid the transfer N times. This layer makes data transfer content-addressed and one-shot:

• a dataset is cut on a FIXED absolute grid of blocks

(BLOB_BLOCK_VALS f64 values = 1 MiB; the last block may be short); each block is keyed by its BLAKE3-256 (stdlib/std/hash_blake3.nu),

• a worker caches blocks on disk ($TMPDIR/swarmb_<hex>.blob), verified

against the hash on arrival — a re-submit or the next iteration round references the same hashes and moves NO data,

• blocks are SEEDED through the ordinary job machinery: a kind_blob

task with the SAME ring key as the compute chunk that needs the block lands on the same worker (consistent hash: same ring + same key → same owner). No new transport protocol, no pull-request re-entrancy — and the cluster HMAC token authenticates blocks exactly like every other payload.

A compute chunk then references its blocks by hash (payload v4, see wasmkernel.nu) and the worker assembles the slice from its cache; a missing or corrupt block fails the chunk VISIBLY (ok=0 → failed_chunks), never silently.

NOTE: siblings are imported by the consumer in dependency order (main.nu: token.nu → blob.nu → …); this file uses token_tag/token_untag from token.nu without importing it, like the other kernel files.

API

@ blob_block_vals → i

The absolute block grid: block b covers dataset values [b·BLOB_BLOCK_VALS, (b+1)·BLOB_BLOCK_VALS).

@ kind_blob → i

Job kind for a block-seed task (kernel=1, wasm=2, wasm_gpu=3).

@ blob_hash ( Vec u ) bytes → ( Vec u )

@ blob_hex ( Vec u ) h → String

@ blob_cached ( Vec u ) hash → b

@ blob_store ( Vec u ) hash ( Vec u ) bytes → b

Verify + store a block (idempotent: an already-cached hash is a no-op). F when the bytes do not hash to hash — a corrupt or forged block is never written under a name it doesn't own.

@ blob_append ( Vec u ) hash ( Vec u ) out → b

Append the cached block hash to out. F on miss or a cache file that no longer matches its hash (deleted / truncated / tampered).

@ blob_seed_payload ( Vec u ) hash ( Vec u ) bytes → ( Vec u )

@ blob_handler ( Vec u ) key → ( @ ( Vec u ) ( Vec u ) )

Worker handler for kind_blob: authenticate, verify content address, cache. Result wire matches the scalar frame ([ok:1][0:8]) so the coordinator's readiness/failure accounting needs no special case.

@ blob_manifest ( Vec u ) data → ( Vec ( Vec u ) )

@ blob_manifest_free ( Vec ( Vec u ) ) m → v

---

wasmkernel.nu

packages/swarm-mcp/src/wasmkernel.nu — phase 2: run a NURL-compiled wasm kernel over a chunk.

The language model writes a kernel as ordinary NURL — anything, not just an expression — and compiles it (via the NURL build service) to a wasm32-wasi module whose main reads lo hi from argv, folds the kernel over [lo, hi), and prints the partial. The coordinator ships that module plus a sub-range to each worker; the worker runs it under the wasm runtime and returns the which the coordinator combines with the reduce op exactly as in phase 1.

The runtime is nwasm, the pure-NURL WebAssembly runtime, compiled INTO this binary (deps/nwasm) and run in-process by default: a worker needs no external runtime, no subprocess and no writable filesystem — which is what lets a unikernel guest run wasm chunks at all. Setting $NURL_WASM_RUNTIME switches CPU chunks to an external binary over the CLI contract <runtime> run <module> <lo> <hi> (partial on stdout) — the external Bytecode-Alliance wasmtime for its JIT, or the nwasm CLI; GPU chunks always use that external contract.

CPU chunk (kind_wasm) : [lo:8 BE][hi:8 BE][wasm…] GPU chunk (kind_wasm_gpu) : payload v2 — see wasm_gpu_chunk_payload below (adds an output mode, K, and runtime params)

The reduce op is baked into the module by the wrapper, so the partial is final per chunk. A scalar module prints its partial as a decimal integer on stdout; for a float task that integer is the partial's f64 BIT PATTERN (see buildwasm.nu), so the same int wire carries it unchanged — the coordinator reinterprets it. Vector modes (sample/hist) write raw little-endian f64s to a worker-named output file instead (fast: one fwrite, no per-value stdout).

The module is cached on each worker by a content hash, so re-running the same kernel (every chunk of a task, and across tasks — runtime params don't touch the module) writes the .wasm only once.

API

@ kind_wasm → i

@ kind_wasm_gpu → i

GPU wasm chunks are a SEPARATE job kind, routed on the GPU capability ring (dist/job job_set_ring): only --gpu workers own these keys, and only they register this handler — a non-GPU worker can never receive (or garble) one.

@ gpu_mode_scalar → i

GPU task output modes (baked into the generated program AND carried in the chunk so the worker knows how to invoke the module and collect its output).

@ gpu_mode_sample → i

@ gpu_mode_hist → i

@ gpu_mode_vecreduce → i

@ gpu_mode_shuffle_map → i

@ gpu_mode_shuffle_reduce → i

@ data_dtype_f64 → i

Dataset storage-type codes (shared by the worker and the generator). A dataset element is stored in its native width and promoted to a double on the GPU. Defined here (the lowest shared module) so both wasmkernel and cudakernel see them. Code 1 = f64 keeps the original single-width behaviour.

@ data_dtype_f32 → i

@ data_dtype_i32 → i

@ data_dtype_i64 → i

@ data_dtype_esz i dtype → i

@ wasm_chunk_payload i lo i hi ( Vec u ) wasm → ( Vec u )

@ wasm_gpu_chunk_payload i mode i lo i hi i kbins ( Vec i ) params ( Vec u ) data ( Vec u ) wasm → ( Vec u )

── GPU chunk payload (v2 / v3) ────────────────────────────────── v2: [2][mode:1][lo:8][hi:8][K:4][nparams:2][f64 bits:8×n][wasm…] v3: [3][mode:1][lo:8][hi:8][K:4][nparams:2][f64 bits:8×n] [dlen:4][data slice: raw LE f64s][wasm…] mode selects the module's output contract; K is the histogram bin count (0 otherwise); params are runtime kernel arguments passed on the module's argv as f64-bit-pattern decimals — the SAME module (same content hash, no rebuild, warm worker cache) serves any parameter values. v3 additionally carries the chunk's DATASET SLICE — data[lo..hi) as raw doubles, following the map-reduce rule that a split travels with its task. dlen is exact (8·(hi−lo)) so a truncated frame is rejected, and dlen=0 means "no data" (the encoder then emits plain v2).

@ wasm_gpu_chunk_payload_blobs i mode i lo i hi i kbins ( Vec i ) params i b0 i dtype ( Vec ( Vec u ) ) hashes ( Vec u ) wasm → ( Vec u )

v4: the chunk references its dataset blocks by CONTENT HASH instead of carrying the slice — [4][mode:1][lo:8][hi:8][K:4][nparams:2][f64 bits:8×n] [b0:4][nblob:2][hash32 × nblob][wasm…] b0 is the absolute index of the first referenced block on the dataset's block grid (blob.nu); the worker assembles data[lo..hi) from its block cache (seeded via kind_blob) and FAILS the chunk on any missing block.

: GpuChunk

: GpuChunk {
    i ok
    i mode
    i lo
    i hi
    i kbins
    ( Vec i ) params
    ( Vec u ) data  // raw LE f64 slice (empty without a dataset)
    i b0  // v4: first referenced block index on the dataset grid
    i dtype  // v4: storage type code (1 f64 · 2 f32 · 3 i32 · 4 i64)
    ( Vec ( Vec u ) ) blobs  // v4: 32-byte block hashes (empty otherwise)
    ( Vec u ) wasm
}

@ gpu_chunk_free GpuChunk c → v

@ wasm_gpu_chunk_decode ( Vec u ) body → GpuChunk

@ gpu_chunk_assemble inout GpuChunk c → b

v4: assemble data[lo..hi) into c.data from the content-addressed block cache. A v2/v3 chunk (no blob references) is a no-op T. F on ANY missing or corrupt block, or a reference window that does not cover the range — the caller fails the chunk VISIBLY instead of computing on garbage.

@ _wasm_hash ( Vec u ) v → String

FNV-1a/64 content hash → 16 lowercase hex chars, for the cache filename.

@ wasm_runtime_probe → String

Can this node actually run wasm modules? A worker used to accept chunks with no runtime installed and only reveal it as an unexplained failed chunk once a task arrived; the node knows the answer at startup, so it should say so then. Returns "" when the runtime works, else the reason it does not.

@ wasm_runtime_gpu_capable → b

Does the resolved runtime understand --allow-gpu? Only the pure-NURL nwasm bridges CUDA/NVRTC host imports, and a --gpu worker without it can never run a GPU chunk. Checked by asking for its help text.

@ chunk_err_marker → i

── why a chunk failed ─────────────────────────────────────────── A failed chunk used to be a bare ok=0: the coordinator could count failures but never say why, and the worker's real reason (no wasm runtime on PATH, a module trap, a CUDA error) was thrown away — leaving "failed_chunks: 2" as the whole story an agent gets. The reason now rides home as a TRAILING suffix on the result frame:

… existing frame … [utf8 message][len:2 BE][0xE7]

Appending at the END keeps every existing frame parse byte-identical (both the scalar [ok][partial] and the vector [ok][count][f64…] readers work off leading offsets), and the 0xE7 marker makes the suffix self-describing, so a coordinator can look for it without knowing which frame shape it is holding. An older worker simply never appends one, and an older coordinator ignores the tail — mixed-version clusters stay sound.

@ chunk_err_max → i

@ chunk_err_push ( Vec u ) r s msg → v

@ chunk_err_read ( Vec u ) body → String

The message a failed chunk carried, or an empty String when it carried none.

: WasmRun { i ok i value String err }

Run a cached module over [lo, hi). ok=1 iff the runtime ran the module to a zero exit — a failed chunk (missing runtime, module trap, GPU error path exiting non-zero) is REPORTED, not folded into the reduce as a silent zero. allow_gpu≠0 passes --allow-gpu so the module's env imports (CUDA/NVRTC) resolve to real hardware — this needs the pure-NURL packages/nwasm.

@ wasm_handler ( Vec u ) key → ( @ ( Vec u ) ( Vec u ) )

The worker handler for a (CPU) wasm chunk: verify the cluster HMAC tag, cache the module by content hash, run it under the wasm runtime, return the partial tagged for the coordinator. An untrusted/forged payload yields a tagged ok=0 — a stranger can't make a worker fetch-and-run an arbitrary module. key is the cluster HMAC key (token_key), captured by the handler closure.

: GpuOut { i ok i scalar ( Vec u ) bytes String err }

@ wasm_gpu_handler ( Vec u ) key → ( @ ( Vec u ) ( Vec u ) )

GPU worker handler: decode payload v2, cache the module, run per mode, and reply the mode's result frame — scalar [ok:1][bits:8], vector [ok:1][count:4 BE][f64 LE × count]. Forged/undecodable payloads → ok=0.