spec: specialized model runtimes

[ndizazzo]

Prototype spec: specialized model runtimes

Status

Draft prototype specification. All decisions required to start have been made.

The prototype builds a MeshLLM-owned inference engine specialized for one closed target:

model identity x weight profile x GPU architecture

The first target is the same target used by NInfer: Qwen3.8-27B on an RTX 5090-class Blackwell GPU. NInfer is the external correctness, quality, and performance control. MeshLLM does not fork, vendor, link, parse, package, or distribute NInfer code or artifacts. Both engines consume independently converted weights from the same upstream BF16 checkpoint.

This is a local and internal prototype. It must not be published to member hardware until the deferred production trust and ABI work is complete.

Goals

• Own the runtime end to end in Rust: artifact format, conversion, execution, kernels, serving ABI, and tests.
• Resolve model architecture, weight layout, execution schedule, and kernel choice at build time for a small declared set of targets.
• Select a specialization only when the requested model artifact and host GPU both match. Fall through to llama.cpp everywhere else.
• Make additional targets repeatable engineering work. Hand-tuning per model and GPU is expected, but the runtime structure, recipe, tests, and recorded knowledge must be reusable.
• Require independent numerical evidence for every specialization.

Non-goals

• A general model graph or a replacement for llama.cpp.
• Runtime plugin discovery or string-driven kernel dispatch.
• CPU offload or multi-GPU execution inside one specialization.
• Competitive performance on an untuned GPU vendor.
• Vision, MTP speculative decode, long context, or batching in the first correct text milestone.
• Public distribution during the prototype.

Architecture

A specialization is a native runtime loaded in-process behind the existing Skippy C ABI. It is not an HTTP plugin or sidecar. MeshLLM already has the right host boundary: a backend-neutral host, hardware-aware runtime resolver, dynamic library loader, and llama.cpp fallback (docs/design/NATIVE_RUNTIMES.md).

T3  Distribution   manifest, model identity, hardware match, rank
T2  Serving        Skippy ABI, sessions, KV state, request lifecycle
T1  Execution      fixed schedule, arena, graph capture, tokenizer, sampler
T0  Kernels        GEMM, attention/GDN, quant/dequant, norms, RoPE

T3 is mostly existing MeshLLM code. T2 and T1 are new Rust. T0 is the principal technical risk and the source of the expected speedup.

Runtime selection

Add a model-identity constraint to NativeRuntimeArtifact:

#[serde(default, skip_serializing_if = "Vec::is_empty")]
pub serves: Vec<ModelIdentity>, // { model_id, weights_id }

Add a matching resolver rejection:

ModelNotServed {
    requested: ModelIdentity,
    serves: Vec<ModelIdentity>,
}

An empty serves list preserves today's general-purpose runtime behavior. Selection must evaluate this conjunction in order:

1. The requested model resolves to a concrete artifact already present on the host. Selection never downloads a specialization as a routing side effect.
2. The runtime declares the exact (model_id, weights_id) in serves.
3. The host satisfies the runtime's GPU architecture, driver, toolkit, and VRAM requirements.
4. The highest-ranked matching runtime wins.
5. No match falls through to llama.cpp unchanged.

Matching a model name without a concrete weight identity is invalid. A GGUF and a specialized artifact can represent the same model name while requiring different loaders and kernels.

Code ownership

The prototype lives under crates/mesh-specialize/ with clear tier boundaries:

crates/mesh-specialize/
  AGENTS.md
  KNOWLEDGE/
  abi/
  engine/
  kernels/
  packages/qwen3_8_27b/
  reference/

Exact crate boundaries can change during implementation. The invariants cannot: vendor-specific code stays in the kernel tier, model-specific execution stays in the package, and reusable runtime behavior stays above both.

Kernel and portability contract

The speed-first kernel decision is final:

No C or C++ source, no CMake, nvcc, or hipcc build path, and no vendored native compute library in the MeshLLM tree. Kernels are authored in Rust. Inline asm! PTX is permitted and expected in hot loops. The vendor driver is a dynamically loaded system dependency.

NInfer is direct evidence for this boundary. Its CUDA implementation owns its compute kernels and still uses inline PTX for tensor-core MMA, ldmatrix, cp.async, and register budgeting. The Blackwell NVFP4 MMA instruction has no C++ intrinsic. Avoiding assembly would remove the instructions needed for the performance target.

Scalability comes from three enforced rules:

1. Keep an independent reference kernel for every op. Write a clear, slow implementation first and never delete it. It is the differential oracle for each tuned kernel and the correctness-first bring-up path for a new GPU.
2. Generate shape variants from recipes. Constants such as tile shape, dtype, head shape, and architecture are generator inputs. Do not copy whole kernels per model when code generation can express the variation.
3. Inventory every assembly site. Each asm! block records its owning op, required architecture, and reference sibling in KNOWLEDGE/asm-inventory.md. No vendor type, symbol, or conditional may appear above T0.

The portability promise is intentionally narrow: a new vendor may require a full set of hand-tuned kernels, but it must not require changes to the artifact format, resolver, ABI, session model, execution structure, or conversion pipeline. A new vendor reaches correctness through the device abstraction and reference kernels, then tunes op by op.

Feasibility gate

The load-bearing unknown is whether Rust's nvptx64-nvidia-cuda path can emit the required instructions with usable register allocation. Test the hardest instructions first:

• mma.sync.aligned.m16n8k64 ... e2m1.e2m1.f32.ue4m3
• ldmatrix.sync.aligned.m8n8.x{2,4}[.trans]
• cp.async.{cg,ca}.shared.global and wait groups
• setmaxnreg.{inc,dec}.sync.aligned.u32
• ordinary BF16, FP16, and INT8 mma.sync shapes

Then implement one quantized GEMM plus fused RMSNorm, with a portable sibling and an independent CUDA-library reference used only by the test harness. That reference is never linked into, shipped with, or called by the runtime. Record correctness, emitted PTX, register pressure, performance, profiler usability, driver/toolkit versions, clock state, and commit.

If Rust cannot emit the tensor-core path correctly, stop before building the model runtime. Continuing would require a new decision that relaxes the Rust kernel constraint.

Target package and artifact

First package

Target Qwen3.8-27B directly. Do not add a small-model milestone.

The target is a demanding hybrid model:

• 64 layers: 16 full-attention and 48 GDN layers
• GQA with partial RoPE on full-attention layers
• recurrent GDN convolution state
• optional MTP and vision paths
• 262,144 native context in the reference package

The first correctness milestone implements the text path, greedy decode, batch 1, and short context. Vision, MTP, long context, and batching follow only after the full text path is correct.

The overridable working default is groupwise-int first, then nvfp4 during performance work. This keeps the first end-to-end correctness result independent of the hardest NVFP4 instruction. The Phase 1 feasibility gate still tests NVFP4 first because its failure decides whether the later speed path exists.

.mspec format

Use a simple placement container owned by MeshLLM:

magic     "MESHSPEC\0", format version
directory identity and objects { name, kind, dtype, shape, layout,
                                 offset, length, sha256 }
payload   aligned tensor blobs

The container carries tensor placement and integrity metadata. Execution semantics, schedules, kernel choice, and supported identities are compiled into the package. Unknown identities fail explicitly.

Conversion and qualification are separate operations:

• Convert: upstream BF16 checkpoint to .mspec, tensor inventory, and recipe hash. CPU-only and architecture-independent.
• Qualify: run parity, quality, and performance evidence on the exact target GPU.

Extend the existing crates/model-package/ job machinery for conversion. Do not build a second packaging system.

Prototype ABI scope

The current Skippy loader requires 90 symbols: 64 skippy_*, 22 mtmd_*, and 4 llama/ggml symbols. It resolves each symbol across a list of libraries, with the last provider winning, after checking skippy_abi_version.

For the prototype, ship a separate, deletable mesh-specialize-abi-stubs library that provides the 26 foreign symbols. Each stub logs its symbol name once and returns the clean failure value already handled by its caller. The engine exports only the ABI it owns. Pin skippy_abi_version to the current version and rebuild when it changes.

The stub library must remain separate from the engine so production capability negotiation can remove it cleanly. A do-nothing engine plus stub library must load before model implementation starts.

Deferred production work:

• split mandatory ABI functions from optional capabilities
• negotiate ABI compatibility instead of exact patch equality
• replace llama.cpp log-derived metadata with structured runtime events
• add specialization selection diagnostics
• enforce artifact signing and attestation
• isolate or automatically demote crash-prone in-process runtimes

The prototype remains local and internal until those items are complete.

Knowledge base requirement

Create the knowledge base in the first implementation commit, before the engine exists:

crates/mesh-specialize/
  AGENTS.md
  KNOWLEDGE/
    README.md
    pitfalls/<slug>.md
    findings/<slug>.md
    optimizations/<slug>.md
    dead-ends/<slug>.md
    ah-ha/<slug>.md
    asm-inventory.md

The path-local AGENTS.md indexes the knowledge base and requires agents to read relevant entries before editing. Every pitfall, finding, optimization, dead end, and useful realization gets one file.

Each entry records:

• category and status, including a link when superseded
• exact model, GPU architecture, driver, toolkit, clock state, and commit
• expected and observed behavior
• reproduction command or linked evidence
• measurement method and before/after values for optimizations
• one durable rule for future work

Dead ends are never deleted. Supersede them with a pointer. Negative claims must state exactly where they were checked. A generated artifact cannot verify itself, and tuned kernels must be checked against an independent reference.

Any change under crates/mesh-specialize/** must update the knowledge base or explain kb: none - <reason> in the PR. A claimed optimization must cite its measured knowledge entry.

Seed entries with the findings already earned:

• PTX escape hatches are required for the Blackwell fast path.
• NVFP4 block-scaled MMA has no source-language intrinsic.
• A single-architecture build is a deliberate tradeoff, not a scalable design.
• NativeRuntimeBackendKind::Other skips meaningful hardware validation and must not be used as a shortcut for a new GPU vendor.
• NInfer's HTTP API provides token sequences, not layer logits.

Correctness and control

Every specialization must have:

1. A conversion recipe and tensor inventory.
2. Independent reference kernels for tuned device kernels.
3. An independent Python model oracle that shares no engine code.
4. Per-layer activation parity and end-to-end token fixtures.
5. NInfer comparison on the same model, weights, GPU, prompt set, context, and concurrency.

NInfer serves three control roles:

Gate	Comparison
Correctness	Greedy token sequence must match on fixed fixtures.
Quality	Run the same task-level evaluation set against both engines.
Efficiency	Compare decode tok/s, TTFT, and peak VRAM on the same host.

NInfer's HTTP API does not expose the layer logits required to localize a divergence. The independent Python oracle provides per-layer diagnosis. NInfer answers whether the externally visible result and measured efficiency agree with an independent implementation.

Control runs record both engine versions, MeshLLM commit, model recipe hash, GPU, driver, toolkit, clock state, prompts, context, concurrency, and raw results. NInfer stays a separately executed binary and .ninfer files are never read.

Additional required tests:

• .mspec round-trip and corruption rejection
• table-driven resolver and llama.cpp fall-through behavior
• randomized kernel differential tests, including odd and boundary shapes
• FFI pointer, length, panic, and error-path tests
• fixed-target performance regression tests
• long-running VRAM leak and determinism checks before production work

Development sequence

Phase 0: prove native-runtime integration

• Add serves and ModelNotServed without changing existing manifest behavior.
• Build the separate 26-symbol stub library.
• Load a do-nothing Rust runtime through the existing loader.
• Prove exact-model selection, wrong-model rejection, wrong-hardware rejection, and llama.cpp fall-through.

Exit: MeshLLM can load a Rust implementation of its native ABI and select it only for an exact resident model artifact on a matching host.

Phase 1: prove the Rust kernel path

• Create crates/mesh-specialize/AGENTS.md and KNOWLEDGE/ first.
• Run the instruction-emission feasibility gate.
• Build and measure the quantized GEMM plus RMSNorm test kernel and its portable sibling.

Exit: every required instruction emits and computes correctly, register pressure is usable, the kernel can be profiled, and all results are reproducible from the knowledge base.

Stop if the tensor-core path cannot be expressed correctly in Rust. Phase 2 depends on this result.

Phase 2: make Qwen3.8-27B correct

The Python oracle can start in parallel with Phases 0 and 1. Model runtime work waits for Phases 0 and 1 to pass.

Use oracle activation injection so only one new subsystem is under test at each rung:

Rung	Engine-owned work	Gate
0	Independent Python oracle from upstream BF16	Reproduce published reference outputs.
1	.mspec writer, reader, and loader	Tensor identity and byte parity.
2	Embedding and one full-attention layer	Per-tensor parity.
3	All full-attention layers	Per-layer parity, GDN injected.
4	GDN layers, with convolution state isolated	Per-layer and state parity.
5	Full BF16 text path, greedy, batch 1	Exact token parity with NInfer.
6	groupwise-int under the same gates	Parity and quality hold.

Exit: on the target 5090 host, MeshLLM resolves and loads the specialization, serves correct greedy completions matching NInfer on the fixture set, and falls through cleanly on all non-matching hosts. Performance is not an acceptance gate for this phase.

GPU runner capacity and pinning are handled outside this spec.

Phase 3: make it fast

Optimize in measured rungs: fused attention/GDN, KV layout, graph-captured decode, NVFP4 tensor-core GEMM, batching, then MTP. Every accepted optimization records a before/after knowledge entry; every abandoned path records a dead-end entry.

Exit:

• Floor: at least 2x llama.cpp decode throughput on the same model and hardware.
• Target: parity with or better than NInfer on decode throughput, with matched context and concurrency.
• Correctness and quality gates continue to pass.

Later scaling and production

After the prototype proves one target:

• add a second model before package assumptions ossify
• add a second GPU vendor using the device abstraction and reference kernels, with no changes above T0
• build distributed qualification only after sandboxing and attestations exist
• complete the deferred ABI, signing, diagnostics, and crash-containment work

Completion criteria

The prototype is successful when all of the following are true:

• MeshLLM selects the specialization only for the exact model artifact and supported host, with unchanged llama.cpp fall-through elsewhere.
• Qwen3.8-27B greedy text output matches NInfer on the agreed fixture set.
• The independent oracle localizes model and kernel divergences per layer.
• Decode throughput is at least 2x llama.cpp on the same 5090, with parity-or-better against NInfer as the target.
• The implementation contains no C/C++ source or build path, and every PTX site has a portable sibling and inventory entry.
• The knowledge base contains the evidence, failed approaches, and measured optimizations needed for another agent to add the next model or GPU target.
• No prototype artifact is published to member hardware.

Source anchors

• MeshLLM native runtime boundary and unresolved signing: docs/design/NATIVE_RUNTIMES.md
• Runtime manifest and resolver: crates/mesh-llm-native-runtime/src/manifest.rs, resolver.rs
• Skippy ABI loader: crates/skippy-ffi/src/lib.rs
• Existing package jobs and correctness harness: crates/model-package/src, crates/skippy-correctness/src
• NInfer kernel and control evidence at the researched revision: README.md, mma.cuh, nvfp4_w4a4_tma.cuh, and serving.md