Skip to content

LLM.md

Latest commit

 

History

History
1536 lines (1318 loc) · 118 KB

LLM.md

File metadata and controls

1536 lines (1318 loc) · 118 KB
Raw

LLM.md — Hanzo Engine

Guidance for AI agents working in this repo. (Symlinked to CLAUDE.md / AGENTS.md / QWEN.md / GEMINI.md — see "Context for All AI Assistants" below.)

What this is

Hanzo Engine is Hanzo AI's native inference engine: a fast, multimodal Rust engine (LLM + vision + audio + speech + image + embeddings) that serves the /v1 inference plane of the Open AI Cloud. One hanzo binary (run, serve, bench, tune) + a Rust crate (hanzo) + a Python package (hanzo).

Provenance and licensing — read before touching LICENSE, NOTICE, or a manifest

Hanzo Engine is a fork, not an independent implementation, and the history it was forked with is retained in full. NOTICE records the project, the history and every vendored source with its licence; tools/upstream-map.sh maps the original crate paths onto ours (the crate rename rewrote commit hashes from 2024-02-27 onward, so SHAs stop matching there even where the content is unchanged). A renamed file is still the original project's file.

The licence posture, and why each piece exists:

  • LICENSE is MIT and names two holders, the original author first, then Hanzo AI, Inc. MIT §2 is satisfied by retaining that copyright line, so the first line is load-bearing. Never delete it. This is not hypothetical: 963f0ff2b (2026-06-01) replaced it with a Hanzo-only line as part of a brand sweep, and 3b993189e restored it 45 days later. Releases cut in that window went out non-compliant. Brand sweeps skip LICENSE, LICENSE-APACHE and NOTICE.
  • LICENSE-APACHE is for vendored GPU kernels. The tree bundles Apache-2.0 CUDA/Metal source with per-file headers intact; a ported file's leading header (copyright, licence, where it was adapted from) stays byte-for-byte, and the comments below it describe our code. Apache-2.0 §4(d) makes NOTICE an actual distribution obligation — unlike MIT, retaining the licence text alone is not enough. NOTICE must ship with any binary or crate we distribute.
  • The tensor framework is our fork, hanzoai/ml (hanzo-ml, hanzo-nn, hanzo-transformers, hanzo-flash-attn{,-v3}, hanzo-metal-kernels), consumed as crates.io registry deps under MIT. Attribution for it lives in that repo's LICENSE files and NOTICE — don't restate them here.
  • Workspace license = "MIT" matches the original. hanzo-router declares MIT OR Apache-2.0; that is deliberate and safe — it is first-party routing code with no forked dependency.

The rule that generalises: license = in a manifest is a claim; the LICENSE files are the fact. A fork cannot relicense, and a find-and-replace across manifests or copyright headers is a licensing change, not a chore.

One name: Hanzo Engine

The tree names only itself. Comments, docs, tests, logs, CLI help and identifiers say what our code does, or the math; a reference value is "the reference value from the published config" or "the recorded reference output (the fixture)", never another engine's name. Comparisons with other engines are not recorded here. What legitimately carries a third-party name, and stays: LICENSE, LICENSE-APACHE, NOTICE and tools/upstream-map.sh; the leading copyright/licence/provenance header of a ported file; dependency names in Cargo.toml; data identifiers (Hugging Face repo IDs, file names and fixture keys the code reads, GGUF metadata keys and arch strings, executable names a guard matches on a host); model and format names (Llama, Qwen, Mistral, GGUF, safetensors, ggml quant types); the OpenAI and Anthropic API names, which are our wire protocol. Rewrite by reading each hit: a blind sed once stripped third-party licence notices and renamed real packages into 404s.

Merging upstream (2026-09-15)

Upstream v0.9.3 merges through the crate rename, not around it. A plain git merge grafts: rename detection cannot pair upstream's core crate directory with hanzo-engine/ once contents diverge, so upstream's tree lands beside ours (1,622 duplicates). What works:

  • Branch upstream-renamed: one commit renames upstream's crates, paths and identifiers onto our layout.
  • Merge it, then for each conflicted path run git merge-file --diff3 with the base blob read from the pre-rename path at the fork point 0a14af4a. The 327 conflicts become ordinary three-way hunks, resolved in slices: models, pipeline core, attention and quant, server, root.
  • tools/merge-audit.sh checks that every upstream-added file landed and is declared, and that our-only subsystems (router, NVFP4, ROCm, Vulkan) survived.
  • Files main renamed by hand are invisible to the pairing and need their own three-way: upstream's server builder is our server.rs, its router builder our router.rs, its Python stub our hanzo.pyi. git diff -M --name-status 0a14af4a main lists them.

Decisions in the merge:

  • Upstream's types where both sides had one: Nvfp4Layer (our vecmat/WMMA kernels and the CUTLASS W4A4 prefill with our activation quantizer; no cuTile on the NVFP4 path), crate::gdn (our ROCm, Vulkan, Metal and CUDA backends inside it), the positions-only rotary API, hybrid-cache checkpoint lanes, speculative decoding (MTP, DFlash, rollback, plus our DSpark and prompt lookup), the CPU attention backend, FlashInfer metadata (GQA groups 1..=8 and 16).
  • Ours kept: GGUF quantized models with ROCm/Vulkan/CUDA decode graphs, ROCm and Vulkan attention fast paths, key-chunked eager attention for long prefill, the router, the license check.
  • hanzo-ml takes the framework's newer APIs (barrier_pool, indexed_gemv, gemv_fused_shared_lhs, quantized embedding, safetensors 0.8), synced from its main plus 35d7ae7c, the rev upstream pins.
  • HTTP requests no longer choose filesystem locations: training adapters go to ~/.cache/hanzo/adapters/{id}/{name}, calibration imatrices to ~/.cache/hanzo/cimatrix/{name}.
  • Dropped: upstream releases/ and CITATION.cff, the GitHub workflows and Dockerfiles main had removed, and per-slot seqlen_offsets (upstream's pool zeroes a slot on allocation). The cutile feature stays optional; it needs CUDA 13.2+, which spark does not have yet.
  • Known loss: Qwen2.5-VL video timing in MRoPE (second_per_grid_ts), to re-add inside upstream's compute_rope_index.

evo serving alone: long prompts (2026-09-15)

Measured in a restored window on evo's serving replica (Qwen3.8-27B Q6_K, 3 runs per point): 53 s to first token at 12.5K prompt tokens, and 22.9 s with four clients at 1.3K. That is what the fleet feels when evo serves alone during a spark window. It is the argument for keeping both hosts up and for the router preferring spark on long prompts.

Claude Code against our stack: what upstream does not fix (2026-09-15)

Measured while porting upstream's Anthropic surface. Of the four things that break Claude Code against this cluster, upstream fixes none:

  • tool_reference blocks (what ToolSearch results are) — 400.
  • document blocks in tool results — 400.
  • MCP image shapes (data + mimeType) — base64 dumped into the prompt as text.
  • Server-side tool_choice — half-fixed: it recognises web_search_* and code_execution_* only, so bash_*, text_editor_* and computer_* still 400 for want of a schema.

Its count_tokens restructuring is the real fix and we took it.

Where each is handled matters. hanzo-engine now renders tool_reference natively, but the pool's replicas on spark and evo do not run hanzo-engine yet, so hanzo-router's rewrite is what actually keeps Claude Code working today and must stay until the engine serves the fleet.

Deliberately not ported, because the chains run to ~22k lines and would tear out our router work: upstream's error-codes, skills, /v1/responses and OpenAI tool-surface commits (ApiError, skill_store, OpenAiToolSurface, StreamOutcomeHandle), the tool-call lifecycle and the forced-decode half of required, and the agentic file-surfacing commits, which live in a shell.rs we do not have.

Cluster toolchain (2026-09-15)

  • spark driver 580 caps CUDA at 13.0, so cuTile (needs 13.2+) could not build. CUDA 13.4.1 now sits beside 13.0; /usr/local/cuda still points at 13.0 because spark's serving stack pins that path. Build with CUDA_HOME=/usr/local/cuda-13.4, run with LD_LIBRARY_PATH=/usr/local/cuda-13.4/compat:...lib64, which is a newer user-mode driver on the old kernel module. No driver upgrade, no reboot.
  • A CUDA build script pins the compiler first, then probes the version with that same path. cudaforge resolves NVCC → PATH → CUDA_HOME, so PATH wins by default and a host with two toolkits compiles with one and gates features on the other. hanzo-quant, hanzo-engine, hanzo-paged-attn and hanzo-flash-attn each carry the same pin, and the four must agree.
  • hanzo-ml carries the newer framework APIs the merge needs (barrier_pool, indexed_gemv, gemv_fused_shared_lhs, quantized embedding, safetensors 0.8, GgmlDType::block_align). Storage locks moved to parking_lot, so storage_and_layout() returns StorageRef, not a std guard.
  • tools/merge-audit.sh gates the merge. It fails unless every upstream-added file is present and compiled, deliberately removed with a reason in tools/merge-audit-removed.txt, or renamed to a path that exists; a removal entry whose file reappears is itself a failure. It also knows that main.rs, include! and #[path] compile a file, not only mod.

Releasing the image (2026-10-08)

  • One image: hanzo.yml declares ghcr.io/hanzoai/engine from Dockerfile, linux/amd64, holding hanzo-server alone (the deployment's command). A release is the version in Cargo.toml and the 18 workspace entries of Cargo.lock, a release: hanzo-engine X commit and a vX tag.
  • Nothing builds on a push or a tag: GitHub Actions is off on hanzoai/engine (gh api repos/hanzoai/engine/actions/permissions answers enabled: false), and cargo fmt --check is red across the workspace on main. The image is asked of the door directly, as org hanzo: POST https://api.hanzo.ai/v1/build with {repo, sha, image: ghcr.io/hanzoai/engine:X, tags: [vX], dockerfile: Dockerfile, context: ., platforms: [linux/amd64]} (bearer hanzo --as hanzo auth token, X-Org-Id: hanzo), then GET /v1/build/<id>. A 429 is the org's concurrent-build limit.
  • The door's 30 minutes include the queue. Its Jobs (2 CPU requested, 8 limit, 16 GiB) share the three 8-vCPU CI nodes with the forge's runners and can wait past the deadline unscheduled; a Job still running at the deadline keeps running and pushes while the row reads failed, so ghcr is the record. A cold build compiled in 16m50s on a contended CI node (1.7.103); 1.7.104 built and pushed in 8 minutes.
  • Three defects stopped every image after 1.7.53, each found by a build that died: the Dockerfile capped cargo at two jobs (1.7.102 lets it use the Job's CPUs); hanzo-zen5/zen5-engine-src was a gitlink no .gitmodules named, and BuildKit's submodule update failed before the Dockerfile was read (1.7.103 drops it); the runtime stage copied hanzo-bench, which has had no binary since b2832f1c5 (1.7.104 ships hanzo-server alone). Dockerfile.cuda and Dockerfile.rocm still build and copy hanzo-bench.
  • Digest: docker-content-digest of ghcr.io/v2/hanzoai/engine/manifests/X with a pull token from ghcr.io/token?scope=repository:hanzoai/engine:pull (basic auth with gh auth token).

Canonical role in the Hanzo model

  • A real implementation repo (native inference) — NOT an SDK, NOT a discovery/wrapper repo. It serves the /v1 API that the cloud SDK family calls.
  • Cloud SDKs (Python flagship hanzoai/python-sdk; @hanzo/ai; Go/Rust/C++/Swift/Kotlin) are the clients. SDK completeness: Python > Rust > C++ > Go.
  • DRY: one impl, one place. Keep engine internals here; link out, never duplicate.
  • Whole-model truth: ~/work/hanzo/SDK-ARCHITECTURE.md.

Brand rules (enforce in every doc and commit)

  • Hanzo is a full AI cloud / inference engine — never an "LLM gateway" or "OpenAI-compatible proxy", and never positioned against a proxy product. Speaking the OpenAI / Anthropic wire formats is a feature, not the identity.
  • HTTP paths are /v1/... only — never an /api/ prefix. The CLI UI's own calls are /ui/v1/* (hanzo-cli/src/ui/mod.rs, webui/src/lib/services/api.ts), web-chat's are /v1/*, and hanzo-router serves no /api/ path. hanzo.yml's v1-only gate greps for a regression; route_registry tests the server routes.
  • Zen is our own model family (zen-eco, zen-ultra, Zen5, …); never surface upstream model names as the brand in public docs.
  • Crisp, modern, developer-first voice; no emoji-spam.

Fast orientation (full detail below)

  • Build / run: see Essential Commands. Agent rules: see Rules for AI Assistants. Symlinks: see Context for All AI Assistants.
  • Core crates: hanzo-engine/ (inference library; canonical hanzo_engine::{infer,embed}, consumed by EVM precompiles 0x0201/0x0202) · hanzo-server/ (binary hanzo-server, released as hanzoai) + hanzo-server-core/src/routes.rs (/v1 HTTP routing) · hanzo-cli/ (binary hanzo-engine) · hanzo-pyo3/ (Python) · hanzo/ (Rust API) · hanzo-quant/ · hanzo-paged-attn/ · hanzo-vision/ · hanzo-audio/ · hanzo-llm-mcp/.

Engineering knowledge base

Hanzo-Specific Components

hanzo-engine/ is a library; the binaries are built by other crates.

  • Library — hanzo_engine::* (the canonical inference API for the Hanzo stack):

    • InferenceEngine trait: fn infer(&self, model_id: &[u8;32], prompt: &[u8]) -> Result<Vec<u8>, EngineError>
    • EmbeddingEngine trait: fn embed(&self, dim: usize, text: &[u8]) -> Result<Vec<f32>, EngineError>
    • Process-wide registry (OnceLock-backed): register_inference_engine, register_embedding_engine, infer, embed
    • The engine-backed implementation (hanzo-engine) (handles HF repos and local paths; derives model_id as SHA-256 of source)
    • Consumers: hanzo-vm precompiles 0x0201 (AI inference) and 0x0202 (AI embedding) — they call hanzo_engine::infer / embed synchronously through the registry
    • NOT the routing/pricing crate. The Hamiltonian-Hidden-Markov MarketMaker lives in hanzo-hmm (~/work/hanzo/net/hanzo-hmm) and prices heterogeneous compute; the EVM precompiles depend on hanzo-engine, not on hanzo-hmm.

  • Binaries (names from cargo metadata):

    • hanzo-engine, package hanzo-cli/: serve, run, tune, bench, from-config, doctor. Links hanzo-server-core in process. install.sh installs it with cargo install --git ... hanzo-cli.
    • hanzo-server, package hanzo-server/: the older model-subcommand server. Release CI packages it as hanzoai (.github/workflows/release.yml), and it logs a deprecation warning naming hanzo-cli.
    • Neither binary is hanzo; that name belongs to the Hanzo CLI (hanzoai/cli), whose hanzo engine serve MODEL runs hanzo-engine serve -m MODEL.

Architecture

Hanzo Engine is a Rust workspace containing:

  • All upstream hanzo workspace members (hanzo-engine, hanzo-server, hanzo, hanzo-llm-mcp, …)
  • hanzo-engine/ — library: canonical Hanzo inference + embedding API
  • Local hanzo-ml fork at ../ml/hanzo-{ml,nn,flash-attn,metal-kernels} overrides upstream's hanzo-ml-* crates via [workspace.dependencies] path overrides

The engine provides comprehensive LLM inference with support for text, multimodal (incl. video), image generation, speech, and embeddings through multiple APIs (Rust, Python, OpenAI HTTP, MCP).

Native ML Pipeline

Three sibling crates live alongside hanzo-engine/ and share the workspace's pinned tensor framework (hanzo-ml 0.9.2-alpha.2, half 2.7.1) — single source of truth:

  • hanzo-federation/ — Vendor-neutral federated-training transport, scheduler, and coordinator. Pure-Rust port of zen/gym/src/gym/distributed with byte-identical wire format (canonical BF16 delta blob + HMAC-SHA256 auth). Public surface: Coordinator::new(lab).serve(addr), Worker::new(cfg).run(step, params, apply, data), TransportClient. Built on the engine workspace's axum 0.8 / reqwest 0.13 / tower-http 0.6.
  • hanzo-quant/ — Pure-Rust quantization: BitDelta (1-bit signs + per-tensor scale, ~32x), DeltaQuant (INT2/4/8 grouped symmetric), DeltaSoup (Byzantine-robust trim-mean aggregation). Uses the workspace hanzo-ml so no version split.
  • hanzo-zen5/ — Zen5 inference adapter. Default ffi feature wraps the vendored zen5-engine C runtime (Metal / CUDA / CPU); native feature swaps in a hanzo-ml DeepSeek V4 Flash scaffold using workspace hanzo-ml / hanzo-nn plus the local hanzo-transformers fork.

Composition with hanzo-engine: hanzo-federation ships the cross-host training fabric; hanzo-quant produces the delta blobs that travel on it; hanzo-zen5 plugs zen5 model variants into the InferenceEngine trait registry alongside the engine-backed implementation. hanzod (in ~/work/hanzo/node) pulls all four via workspace paths (hanzo-federation = { path = "../engine/hanzo-federation" }, etc.).

Essential Commands

Building

The root Cargo.toml is a virtual workspace, so every build names a package.

# Check the library compiles
cargo check -p hanzo-engine --no-default-features --features metal

# CLI binary `hanzo-engine` (what install.sh installs)
cargo install --locked --path hanzo-cli --features metal            # macOS
cargo install --locked --path hanzo-cli --features cuda-full        # Linux, CUDA + flash-attn + cuDNN
cargo install --locked --path hanzo-cli --features rocm             # Linux, AMD

# Release server binary target/release/hanzo-server; release.yml packages it as `hanzoai`
cargo build --release -p hanzo-server --no-default-features --features metal

ROCm on AMD APUs (Strix Halo / gfx1151, unified memory)

APUs expose a tiny dedicated-VRAM carve-out (HIP reports ~1 GB) alongside the large unified GTT pool (~the whole system RAM). A raw hipMalloc targets the 1 GB carve-out and OOMs immediately on any real model. Run with unified memory:

HSA_XNACK=1 LD_LIBRARY_PATH=/opt/rocm/lib \
  hanzo-server --port 8080 gguf -m <repo> -f <file.gguf>

HSA_XNACK=1 lets managed allocations page into GTT so the full unified pool is usable. (The build already matches the fused qk-norm-rope cos/sin cache to the activation dtype; without that fix every GGUF MoE decode step errors on gfx1151.)

Testing & Quality

# Run core tests
cargo test -p hanzo-engine -p hanzo-quant -p hanzo-vision

# Format code (uses rustfmt, ruff, clang-format)
make fmt

# Check formatting
cargo fmt --all -- --check

# Run clippy
cargo clippy --workspace --tests --examples -- -D warnings

Running Models

# Release binary. No default port; --serve-ip defaults to 0.0.0.0.
hanzoai --serve-ip 127.0.0.1 --port 1234 run -m <model_id>
hanzoai --serve-ip 127.0.0.1 --port 1234 multi-model --config models.json   # several models, one port
hanzoai --port 1234 gguf -m <repo> -f <file.gguf>
hanzoai -i plain -m <model_id>                                              # interactive, no server

# CLI binary. serve defaults to 0.0.0.0:1234 and advertises over mDNS unless --no-advertise.
hanzo-engine serve -m <model_id>
hanzo-engine run -m <model_id> -i "prompt"

Models

When integrating a new model, make sure it respects all of the varbuilder .pp calls. In Hanzo, a VarBuilder maintains an internal path vector that acts like a “current working directory” for model weights; every call to pp("sub") (alias for push_prefix) clones the builder and appends sub, so successive calls accumulate a dotted prefix such as transformer.h.0 while leaving the original builder untouched . When you eventually call get(...), Hanzo joins that prefix with the tensor name (prefix + "." + name) and looks it up in the checkpoint backend, producing keys that exactly match the dot-separated names emitted by PyTorch’s state_dict/named_parameters, which means PyTorch-trained weights can be loaded without any renaming . This lets you recreate the PyTorch module tree in Rust by “walking” it: e.g. vb.pp("word_embeddings") grabs word_embeddings., while a chain like vb.pp("encoder").pp("layers").pp(i.to_string()) targets keys such as encoder.layers.0., exactly as shown in community tutorials porting Transformers models to Hanzo . As one maintainer put it, the prefix system lets you “cd” around the parameter hierarchy, giving a lightweight namespace mechanism that keeps Hanzo fully compatible with PyTorch naming conventions while remaining ergonomic to use.

You should also look for a model.safetensors.index.json file for the model at hand to verify correct structure.

Architecture Overview

Workspace Structure

Hanzo-Specific

  • hanzo-engine/ - Hanzo's custom inference server and CLI
    • Custom CLI with model management commands
    • OpenAI-compatible HTTP server (port 36900)
    • Status: Compiles successfully with Metal backend (macOS)

Core Components

  • hanzo-engine/ - Core inference engine, model implementations, pipelines
  • hanzo-server/ - CLI binary entry point
  • hanzo-server-core/ - HTTP server routing, OpenAI API implementation
  • hanzo-pyo3/ - Python bindings (PyO3)
  • hanzo/ - High-level Rust API
  • hanzo-vision/ - Vision model support
  • hanzo-quant/ - Quantization implementations (ISQ, GGUF, GPTQ, etc.)
  • hanzo-paged-attn/ - PagedAttention implementation
  • hanzo-audio/ - Audio processing
  • hanzo-llm-mcp/ - Model Context Protocol client
  • hanzo-bench/ - Benchmarking tools

Key Design Patterns

  1. Pipeline Architecture: All models implement the Pipeline trait in hanzo-engine/src/pipeline/mod.rs. Different model types (Plain, GGUF, GGML, Vision) have their own pipeline implementations.

  2. Model Loading: Models are loaded through Loader traits that handle different formats and quantizations. See hanzo-engine/src/loader.rs.

  3. Request Handling: The server uses message passing with Hanzo struct managing a background thread pool. Requests flow through hanzo-engine/src/engine/mod.rs.

  4. Device Management: Automatic and manual device mapping for multi-GPU setups handled in hanzo-engine/src/device_map.rs.

  5. Decode Graphs (CUDA/ROCm): Steady-state single-token decode replays a captured GPU graph instead of relaunching hundreds of kernels/token (pipeline/cuda_graph.rs, pipeline/rocm_graph.rs). Wired for safetensors in NormalPipeline and for GGUF in GGUFPipeline (pipeline/gguf.rs). Toggle with CUDA_GRAPHS=0 / ROCM_GRAPHS=0. ELIGIBILITY INVARIANT (model_supports_decode_graph): only variants whose decode RoPE reads the device metadata.rope_positions buffer (Qwen3, Qwen3MoE) are graph-safe. Host-offset RoPE (RotaryEmbedding::forward -> selected_rope_cache -> cos.narrow, used by Llama/Phi2/Qwen2) and per-forward mRoPE built via Tensor::from_vec (Qwen35 compute_text_mrope) FREEZE the position at capture and emit garbage on replay. GB10 decode wins: 0.6B +32%, 8B +7%, 30B-A3B MoE +25%, all byte-identical to eager.

Adding New Features

When adding new model architectures:

  1. Implement the model in hanzo-engine/src/models/
  2. Add pipeline support in hanzo-engine/src/pipeline/
  3. Update model detection in hanzo-engine/src/pipeline/normal.rs
  4. Add architecture enum variant in hanzo-engine/src/lib.rs
  5. Update CLI args in hanzo-server/src/main.rs

When adding new quantization methods:

  1. Implement in hanzo-quant/src/
  2. Add to quantization loading logic in pipelines
  3. Update documentation in docs/QUANTIZATION.md

Facial Animation (MuseTalk)

hanzo-engine serve --arch musetalk -m <bundle> loads a MuseTalk lip-sync/avatar model and serves /v1/animate (alias /v1/video/lipsync): POST {model, visual, audio, fps?}, audio + visual in, dubbed mp4 out. A still image -> Portrait, a video -> Footage (lip-sync); the pipeline picks the animator via accepts() and the handler muxes the rendered frames + driving audio with ffmpeg.

-m is a self-contained MuseTalk bundle (a local dir or HF repo). All six sources resolve as fixed safetensors sub-paths of the single model_id:

musetalkV15/musetalk.json
musetalkV15/unet.safetensors
sd-vae-ft-mse/config.json
sd-vae-ft-mse/diffusion_pytorch_model.safetensors
whisper/tiny.safetensors
s3fd.safetensors

Selection path: CLI --arch musetalk -> ModelType::Animation -> ModelSelected::Animation -> AnimationLoader (AnimationLoaderType::MuseTalk) in pipeline/animation.rs. Runs on CPU today (Device::Cpu); GPU is a separate effort. A .pth/.pt pickle still loads (load_vb) but safetensors is canonical.

World Model (Oasis)

hanzo-engine/src/diffusion_models/oasis/ is a native port of Etched/Decart Oasis-500M (MIT, code + weights Etched/oasis-500m): an action-conditioned frame-autoregressive latent-diffusion Minecraft world model (a playable generative game engine).

  • vae.rs — ViT-VAE (vit-l-20): 20x20 conv patchify of a 360x640 frame -> ViT blocks -> 16-dim latent [16,18,32]. Attention is 2D pixel-axial RoPE; MLP is exact GELU.
  • dit.rs — DiT-S/2 spatiotemporal DiT (16 blocks): adaLN spatial-attn + causal-temporal-attn, action conditioning via external_cond added to the timestep embedding. tanh-GELU MLPs.
  • rope.rs — interleaved (lucidrains adjacent-pair, is_gptx=false) axial RoPE. Spatial/VAE use pixel freqs, temporal uses lang freqs; freqs verified against the checkpoint.
  • sampling.rs — diffusion-forcing rollout (sigmoid-beta schedule, per-frame noise, sliding window at MAX_FRAMES=32). Fully-hallucinated: only prompt frames are VAE-encoded, all subsequent frames are generated purely in latent space (no encoder in the loop).
  • mod.rs — WorldModel::{load,encode_frames,generate,decode_frames} + 25-key ACTION_KEYS.

Parity vs the torch reference (CPU f32, parity.rs, weight-gated on OASIS_DIR): VAE encode latent cosine 1.000000, decode PSNR 59.73 dB, DiT single-step cosine 1.000000.

Gotcha: hanzo-ml's Linear only matmuls up to rank 4; the DiT keeps rank-5 [B,T,H,W,D] tensors, so all its linears go through apply_linear (flatten leading dims -> matmul -> reshape).

Example: cargo run --release --example oasis_generate -- --frames 32 --steps 10 --out <dir> (prompt image + action stream -> PNG frames; assemble with ffmpeg).

Important Files to Know

  • hanzo-engine/src/engine/mod.rs - Main engine orchestration
  • hanzo-engine/src/pipeline/mod.rs - Pipeline trait and common logic
  • hanzo-server-core/src/routes.rs - HTTP API endpoints
  • hanzo-pyo3/src/lib.rs - Python API entry point
  • hanzo/examples/ - Usage examples for Rust API

Testing Approach

You should always run cargo check/cargo c before returning to make sure code compiles. If code does not compile, only make edits.

Avoid returning TODOs.

  • Unit tests are colocated with source files
  • Integration tests in tests/ directories
  • Use cargo test -p <crate> to test specific components
  • Python tests require building and installing the package first

Common Pitfalls

  1. Feature Flags: Many features are gated behind Cargo features. Always check what features are needed for your use case.
  2. Device Indices: CUDA device selection uses 0-based indexing
  3. Chat Templates: Models may need specific chat templates - check chat_templates/ directory
  4. Quantization: Different quantization methods have different hardware requirements

Qwen3.5/3.6 hybrid (GDN + MoE) and the portable kernel DSL

  • quantized_qwen3_5_moe is the Qwen3.5/3.6 (qwen3_5_moe arch) hybrid model: Gated-DeltaNet (GDN) linear-attention layers interleaved with MoE. Registered end-to-end (GGUF arch Qwen35/Qwen35MoE in pipeline/gguf.rs, Qwen3_5MoeLoader, vision + text variants). GDN lives in hanzo-engine/src/models/gdn.rs.
  • Fused GDN kernel: gdn.rs routes the gated-delta-rule recurrence through ONE fused per-backend kernel (fused_gdn_gating_{cuda,metal} + the fused scan) that drives Qwen3.5/3.6 to near-parity; the portable f32 ops-composed scan (gdn_step_scalar) is the reference and the CPU/Vulkan fallback. The fused-vs-portable A/B env knob was removed; bit-exact CPU gates (gdn_step_matches_reference_seq1, gdn_recurrence_cpu_shapes) are colocated in gdn.rs and compare the two paths directly.
  • hanzo-kernel (the portable kernel DSL, published 0.2.15) is a CODEGEN: one #[kernel(targets(...))] Rust source lowers to CUDA/ROCm/Vulkan/Metal/WebGPU/CPU. Its op library (gdn, norm, rope, attn incl. sdpa_runtime, quant) and the fuse auto-fusion pass (fusion == composition of Map morphisms, one fused_interp launch, zero intermediates) are all bit-exact on CPU. sdpa/sdpa_runtime online-softmax is the structural cure for the 8B flash-collapse. See ml/LLM.md for the DSL and env-var details.

zen6 (Qwen3.8-Flash-Next) — native and GGUF

  • The name is zen6 (upstream: qwen4_exp): module, loader, --arch zen6, GGUF arch, fixtures, scripts. Upstream spellings are read in two places and never written: Qwen4ExpForConditionalGeneration in NormalLoaderType::from_causal_lm_name, and GGUF general.architecture = qwen4exp, a parse alias on GGUFArchitecture::Zen6 whose qwen4exp.* keys Content::from_readers moves under zen6.*.
  • Two doors, one model. The native safetensors checkpoint (Zen6Loader, pipeline/loaders/normal_loaders.rs) and the same tower quantized upstream to GGUF (pipeline/gguf.rs) both reach models/quantized_zen6::ModelWeights, which is the ONLY place the layers are built.
  • The manifest is the contract. models/zen6::manifest derives every tensor the text model reads — name, stored dtype, shape, and a Role of Device or Host — from the config alone. The checkpoint's headers must equal it (apart from the DEFERRED prefixes), so a silent shape or dtype drift in a release checkpoint is a load failure naming the tensor, not a wrong logit. manifest_is_closed_under_config proves the derivation is total, loader_estimate_equals_allocation that the loader reserves exactly what the manifest sums, and deferrals_match_their_prefixes that what is skipped is skipped on purpose.
  • What is where. Routed experts stay PACKED NVFP4 (E2M1 codes, one E4M3 scale per 16, an F32 global scale) and are read straight into the MoE kernel — nothing dequantizes 71 GB to walk it. Dense side layers are block FP8 (E4M3, one F32 scale per 128x128). The rest is BF16. Layer 1's PLE n-gram table is Role::Host: it is hashed and looked up in an mmap'd buffer, never uploaded, and never charged to the KV budget (table_never_charged, hybrid_kv_charges_attention_and_index_only — only the full-attention layers and their QSA index entries are).
  • W4A4, matched to the server, not to an ideal. hanzo-quant/src/quantize/ reproduces the served activation quantizers bit-exactly — per-token/per-128-group FP8 in the three forms the compiled graph uses, and per-16 NVFP4 with the served fast-math formula. On CUDA that means the approximate instructions (div.full.f32, rcp.approx.ftz.f32) as inline PTX in kernels/quantize, built WITHOUT fast-math so only the places the server approximates are approximate. On CPU the same formulas use IEEE divide/reciprocal and differ in the last bit — which is why the CPU golden is a transcriptor and the GPU vectors are the arbiter.
  • The reference is the server, not the module source. The CPU golden, hanzo-engine/tests/fixtures/zen6.safetensors, holds 24 tokens through the embedding, layers 0-3 (three gated-delta-nets, then attention), the final hyper-connection mixer and lm_head rows [0, 8192), every intermediate recorded. It rounds to bf16 ONLY where a served kernel stores (tl.store or a custom kernel's output); between store points arithmetic is f32, because the compiler fuses away the eager casts. The served-kernel vectors (hanzo-engine/tests/fixtures/zen6_vectors.safetensors, and hanzo-quant/tests/fixtures/vllm_vectors.safetensors for the quantizers) are the served kernels run on the same inputs, so they measure the golden's gap. The recorders drive the reference server and are not in this tree; the fixtures are what the tests read.
  • Two facts about the checkpoint and the recorder. Layer 1's n-gram table is 128 shards in one contiguous run, but the writer put 126 of them in numeric order and transposed one pair, so shard 100 and 101 sit between 97 and 98 — read each shard from its own header range, never from a stride. And six served prefill vectors (gdn.prefill.{conv,q,k,v,core,normed}) are every zero: the recorder's causal_conv1d_fn call returned its empty_like buffer unwritten, so those stages have no served reference at all and their *.golden gaps in the vectors metadata are fiction. g and beta are right, because they come from a and b, not from conv. vectors() refuses an all-zero stage by name. Re-recording those stages settles the call form.
  • Deferred. mtp. (M3: multi-token prediction) is not read by the native door; the GGUF door drafts with the MTP head GGUF published beside the model (models/quantized_zen6_mtp.rs).

QSA sparse attention and YaRN (both doors)

  • Context: the checkpoint's own (the GGUF's context_length, the config's max_position_embeddings, 262,144 either way), times the YaRN factor when one is declared. Yarn::mrope builds it for both doors from what each checkpoint says: the GGUF {arch}.rope.scaling.{type, factor, original_context_length, yarn_beta_fast, yarn_beta_slow} metadata, or rope_parameters with rope_type: "yarn", factor, original_max_position_embeddings, beta_fast, beta_slow (the original defaults to the context either way). The released GGUF carries none, so --rope-yarn-factor F (GGUF serve/run) writes those keys into the loaded metadata (Content::extend_context) before any model reads them: --rope-yarn-factor 4 = the card's static YaRN, 1,048,576 positions. The paged KV cache does not follow the factor: size it with --pa-context-len N (or --max-seq-len, the device-map demand).
  • YaRN on M-RoPE (layers.rs Yarn): the M-RoPE table spans factor times the original length and YaRN counts its rotations over THAT, so the ramp is frequencies 16..24 of 32, not the 14..22 a count over the original length gives. The attention factor 0.1 ln F + 1 multiplies cos and sin (f32, before the cast), so the indexer's roped dims carry it too.
  • QSA (models/qsa.rs): per attention layer an index cache entry of 1 head x 128, K = compressed key at each group's last token, V = raw key per token. Paged: entries 48 + l in kv_layers() (IndexEntries, which ContentConfig and the native door's KvLayers::indexed share), read by layer l. Eager: HybridCacheConfig::side_kv caches after the n-gram pool, batched/split/trimmed with the attention caches. Within the width (2051 tokens) the layer runs its dense attention unchanged; past it: device scores, host DET top-k, union-of-selections masked attention with the QSA kernel's numerics (qsa::softmax_pv, shared with the dense qsa::attend). One forward (Layer::forward -> qsa::forward) serves both doors; only the loaders differ.
  • The two doors' indexers are the same tensors. The safetensors checkpoint stores one self_attn.indexer.index_qk_proj [(4 + 1) x 128, 2560] bf16, the query heads' rows then the key head's, and Gemma norms q_layernorm/k_layernorm as w; the GGUF splits the projection into indexer.q_proj/k_proj and stores the norms as 1 + w in f32. Checked on layer 3 of the released files: rows 0..512 are byte-equal to q_proj, rows 512..640 to k_proj, and the GGUF norms are exactly 1 + w. qsa::Indexer::new loads the former, Indexer::from_gguf the latter.
  • Validate on a GPU by comparing paged vs eager (--paged-attn off) on the same long prompt: the paged path's only code not run by CPU tests is the reshape_and_cache write of the index entry.

Kernel tuning + new-model onboarding: the roofline-driven process (SOTA loop)

A repeatable, measurement-first process. The scar tissue behind every step is in ml/LLM.md ("Vulkan PREFILL: the Q6_K lever, the roofline method..."). The rule: measure in-engine, never trust a cache-warm microbench for a memory-bound kernel, and kill a lever the moment the engine says it lost.

Tune a kernel (the loop that took prefill 212 -> 704 t/s)

  1. PROFILE a real forward: VK_PROFILE_GPU=1 VK_PROFILE=1 hanzo-engine bench --prompt-len 512 --gen-len 0 ... -> per-op GPU time + record/submit/fence + pool fresh/hit counters.
  2. RANK by roofline: achieved BW = bytes-moved / GPU-time vs device peak. Biggest (time-share x distance-from-roofline) = the target. (Q4_K GEMM: ~24 GB/s vs ~135 practical = the 18% that IS the gap.)
  3. PICK THE LEVER by what the roofline says, not by fashion: coalescing / occupancy / layout / reuse / fusion / vectorization. NOT tensor-core dtype (coopmat is measured-dead on gfx1151), NOT bigger tiles (occupancy collapse, measured 2.4x WORSE).
  4. IMPLEMENT in the DSL/shader; GATE BIT-EXACT -- a CPU oracle OR a CONTROLLED known-answer test (uniform weight -> exact k localized the Q6_K decode bug in minutes where random-vs-oracle only said "wrong").
  5. MEASURE IN-ENGINE (--gen-len 0 = the true cold-weight prefill floor). Env-gate A/B variants (VK_Q4K_BM, VK_Q6K_LEGACY, VK_Q4K_COOPMAT). Cache-warm microbench LIES for memory-bound kernels.
  6. KEEP only if it wins in-engine; RECORD the dead ends (this session killed coopmat + bigger-BM).

Autotune (mechanical, per shape/device -- NOT AutoML)

Kernel-level LATENCY search over tile/occupancy configs (BM/BN/BK/threads/unroll), measured and cached per (kernel, shape, device). Deterministic objective (ns), cheap trials -> a for-loop + cache, like triton.autotune / cubecl autotune. The env A/B knobs (VK_Q4K_BM=64/128/256) are the manual version; the DSL autotune automates the sweep so a hand-picked constant can't be a local optimum on another GPU.

Onboard a new model fast (get it to SOTA on our stack)

  1. LOAD the GGUF; confirm COHERENCE with a known-answer generation (hanzo-engine run -m <model> -i "The capital of France is").
  2. PROFILE prefill + decode -> the hot kernels for THIS model's quant mix + shapes.
  3. ROOFLINE-RANK -> the offenders. The classic miss: a quant type with NO tiled prefill path (Q6_K was 76% of prefill because attn_v/ffn_down are Q6_K in Q4_K_M and only had the column matvec).
  4. ENSURE every quant type in the model has a tiled prefill path (Q4_K, Q6_K done; Q5_K/Q8_0/IQ next).
  5. AUTOTUNE the tile configs for the model's shapes.
  6. GATE on the box's roofline: hanzo-engine bench --prompt-len 512 --gen-len 128 auto -m <dir> -f <gguf>; prefill T/s against the compute roof, decode T/s against bandwidth / bytes-per-token. That fraction is the SOTA gate.

The tooling

  • Roofline dashboard (visual "find" UI: roofline plot + priority-ranked kernels + tuning log).
  • VK_PROFILE_GPU (per-op GPU ns), VK_PROFILE (record/submit/fence + pool counters), VK_ROOFLINE (planned: per-dispatch exact bytes/FLOPs joined to timestamps, auto-refreshing the dashboard).
  • The DSL end state (ml/LLM.md "DSL MAGIC"): layout-as-value + a coalesced layout-aware loader + tiled-matmul + pluggable quant-decode + autotune, so every kernel inherits coalescing correct-by- construction and naive strided reads stop being possible.

Qwen3-VL GGUF (text backbone) — qwen3vl / qwen3vlmoe

  • GGUF fast-path for Qwen3-VL: the dense text backbone (qwen3vl, 2/4/8/32B) reuses quantized_qwen3::ModelWeights; the MoE text backbone (qwen3vlmoe, 30B-A3B / 235B-A22B) reuses quantized_qwen3_moe::ModelWeights. Wired as GGUFArchitecture::Qwen3Vl / Qwen3VlMoE in gguf/mod.rs, dispatched in pipeline/gguf.rs, sized in utils/gguf_metadata.rs. No new modeling code — the VL text tower is structurally identical to Qwen3/Qwen3MoE (same blk.* tensors, q/k-norm, GQA); metadata is read under the file's own qwen3vl.* / qwen3vlmoe.* prefix (dynamic path_prefix), so rope.freq_base, head dims and expert counts come straight from the file.
  • Why reuse is bit-correct for text: Qwen3-VL uses interleaved-MRoPE (sections [24,20,20]), but for text-only tokens t==h==w, so every frequency band receives the same position and the interleaved partition is a no-op — identical to the standard 1D RoPE already in quantized_qwen3. Verified: Qwen3-VL-8B-Instruct Q4_K_M loads (36 layers, 151936-tok Qwen3 BPE), dispatches to Model::Qwen3, and decodes coherent text.
  • Image-blind (mmproj not wired): GGUF Qwen3-VL ships TWO files — the LLM gguf (loaded here) and a separate mmproj-*.gguf vision tower (general.architecture=clip, clip.projector_type=qwen3vl_merger, v.blk.* / mm.* / v.deepstack.*, DeepStack at vision layers 8/16/24). GGUFPipeline is text-only (modalities Text→Text, no image inputs processor), so the mmproj is NOT consumed. Full image support needs a clip/mmproj GGUF loader feeding the existing vision_models/qwen3_vl tower PLUS a multimodal quantized-VL text forward (3D-MRoPE position_ids + DeepStack embed injection + <|image_pad|> merge) — a new subsystem, not a reuse of quantized_qwen3::forward (whose signature carries only 1D start_offsets and no DeepStack hook). See the safetensors Qwen3VLLoader for the reference multimodal path.
  • Decode graphs EXCLUDED for qwen3vl/qwen3vlmoe GGUF: they dispatch to the Qwen3/Qwen3MoE model variants, which are already graph-eligible for pure-text 1D RoPE; the VL archs run the same eager/graph text path. (Multimodal 3D-MRoPE, if added, would be position-dependent and must stay eager per model_supports_decode_graph.)

Environment variables (bare names, one-way; de-brand DONE)

The env-flag convention is BARE names (no HANZO_ brand prefix). The de-brand is COMPLETE: every runtime knob is a bare name and the one-off dev A/B "fallback" toggles are DELETED (production always runs the fast path; runtime HW auto-select is kept). One bare name per real knob. Canonical flags in perf_flags.rs:

  • CUDA_GRAPHS, ROCM_GRAPHS, METAL_GRAPHS, FLASHINFER_DECODE, FLASH_PREFILL -- all default ON, set =0 to force the eager/unfused path. (HANZO_ROCM_FLASH_ATTN and its A/B toggle were already deleted; ROCm flash is always-on when applicable.) FLASH_PREFILL gates using_flash_attn() (CUDA).
  • De-branded runtime config: KV_SPILL_DIR, KV_SPILL_BUDGET_MB, ISQ_SINGLETHREAD, VK_FUSED_QKNORM, MXFP4_DP4A, MN_LOCAL_WORLD_SIZE, NO_NCCL, FFI_MODELS, FFI_TOK_DIR, CUDA_FLASH_BF16, METAL_PRECOMPILE, ROCM_GFX_ARCH.
  • DELETED (one-off dev A/B toggles that forced the OLD/slow path): HANZO_GDN_FUSED_FALLBACK, HANZO_ADD_RMSNORM_FALLBACK, HANZO_QK_NORM_ROPE_FALLBACK, HANZO_NO_MEMPOOL_FIX, SAMPLER_TRACE (plus the ml HANZO_Q*K_FALLBACK / HANZO_MOE_*_FALLBACK / HANZO_IQ*_FALLBACK family). The fused vs unfused decision is now the always-fast path; bit-exact oracle tests force the unfused/scalar leg via test-only programmatic setters (layers::set_force_unfused_qk_norm_rope, hanzo_ml::set_force_scalar_matvec), never env, never set in production.
  • License gate reads HANZO_ENGINE_LICENSE_TOKEN / HANZO_ENGINE_LICENSE_FILE (see license.rs); these are product-namespaced identity vars, not perf flags, and stay branded by design -- the ONLY remaining HANZO_ env names.

License key

  • license.rs HANZO_LICENSE_PUBKEY is the public half of the production seed in Hanzo KMS at /engine/LICENSE_SIGNING_KEY (org hanzo, env prod). It is the only key the verifier trusts, in every profile; there is no dev key outside #[cfg(test)].
  • The engine never signs. hanzoai/licensing mints (KMSSigner, same public key as its Pubkey), and hanzo-inc/cloud serves it at /v1/licensing.
  • The tests pin two tokens: one the production key signed (app "test", expired) and one the retired development key signed, which must fail BadSig. Rotating the key means a new KMS value, both public constants, and a fresh production token here and in hanzoai/licensing.

Latest Upstream Features (as of commit 530463af1)

  • Qwen 3 VL - Vision-language model support (#1657)
  • Paged Attention Refactor - Simplified paged attention modules (#1654)
  • Audio Processing - normalize, apply_fade, remove_dc_offset functions (#1572)
  • Gemma 3N - Support for cases where q != (k=v) devices (#1653)
  • No Busyloop Refactor - Improved engine efficiency (#1655)

See docs/ directory for detailed documentation on specific models and features.

Embeddings

  • Path: /v1/embeddings (hanzo-server-core embeddings.rs) sends each input as its own sequence; the scheduler buckets sequences by exact length, so an embedding batch never holds padding and pooling reads every position (embedding_models/layers.rs). The model's token states pass through the sentence-transformers modules of modules.json (Pooling, Dense, Normalize).
  • Pooling reads both 1_Pooling/config.json layouts: sentence-transformers 6 (embedding_dimension, pooling_mode a mode or a list) and the older pooling_mode_* flags. An unsupported setting (include_prompt: false, no mode) fails at load, not at the first request.
  • dimensions keeps the leading values and L2-normalizes them again (Matryoshka); 422 above the model's width. Meaningful for EmbeddingGemma 2 (768/512/256/128) and Qwen3 Embedding.
  • Revision: a multi-model config entry takes revision (hanzo-server's JSON, server-core ModelConfig::revision), passed to load_model_from_hf; a commit pins the weights an index was built from. Single-model mode and the CLI's TOML config load main.
  • EmbeddingGemma 2 (embedding_models/embedding_gemma2.rs, arch embeddinggemma2, auto from EmbeddingGemma2Model): the text backbone only (language_model.*; the vision and audio towers are never read). Gemma 4 layer: norms scale by the weight (not 1 + weight), v is RMS-normed with no scale, attention scale 1, full layers take per_layer_config's head dim 512 and one KV head, sliding layers see |i - j| <= sliding_window (inclusive: the exclusive masker gets window + 1), a projection-only PLE (each layer reads its 512 rows of ple.per_layer_model_projection, shared norm), then layer_scalar, final norm and embedding_projection 512 -> 768 per token. Context 8,192 (CONTEXT, the model card's). Never F16: its activations overflow it. Parity (2026-10-08, evo CPU F32, 64 texts in 14 languages incl. 4 of 250-1,256 tokens, vs sentence-transformers 6.1.0 / transformers 5.19.0): min cosine 0.9999993 at 768/512/256/128, token counts 64/64 equal. Qwen3-Embedding-0.6B on the same build: 0.9999956.
  • Production (universe charts/app/values/hanzo/engine.yaml, 1.7.104): both models in one multi-model process per replica, files/engine/models.json pinning each revision; zen-embedding is Qwen3, zen-embedding-2 is EmbeddingGemma 2. kubectl top 4.9-5.0 GiB a pod (3.2 GiB with Qwen3 alone).
  • Cost on 4 Zen 5 cores (evo cpus 4-7, F32, SciFact passages ~310 tokens, one request at a time): EmbeddingGemma 2 1.26 texts/s (385 tok/s), 1.66 GiB loaded; Qwen3-Embedding-0.6B 0.68 texts/s (222 tok/s), 3.13 GiB; both in one multi-model process 4.69 GiB loaded, 4.93 GiB peak, each at its own speed. torch on the same cores reads 2.8 and 0.9 texts/s. Profile (EmbeddingGemma 2): gemm's AVX2 microkernel 35% (no AVX-512 kernel is selected), hanzo-ml's BarrierPool spin 28%, and MatMul::matmul casts CPU attention to F16 (hanzo-quant lib.rs).

Syncing with Upstream

To pull latest changes from upstream:

# Fetch upstream changes
git fetch upstream

# View what's new
git log HEAD..upstream/master --oneline

# Merge upstream changes (creates merge commit)
git merge upstream/master

# Or rebase Hanzo changes on top of upstream
git rebase upstream/master

# After resolving conflicts, test build
cargo check --package hanzo-engine --no-default-features --features metal

hanzo-router -- the routing mechanism

hanzo-router is the routing MECHANISM (registry, SLO gate, placement, dispatch). It exposes one seam: the RoutePolicy trait -- route(&Request, &User, &Slo, &Registry) -> Route { model, level, modality, confidence }. The rule-based Policy implements it as the cold-start fallback (placement-agnostic, fixed low confidence). New vocabulary lives in registry (Modality, Level) and route (Route, User, Slo, REFUSED_MODEL). Pre-existing API unchanged.

A learned policy implements RoutePolicy outside this repo; its fitted heads mount with ROUTER_HEADS=<path> (Policy::load_heads).

Observability: /health, /metrics, /cache, the request ledger

The bodies are built in hanzo-server-core/src/observe.rs from plain values, the handlers in handlers.rs.

  • Nothing waits on the engine. The engine loop publishes to its IntervalLogger: a heartbeat (Phase idle/loop/prefill/decode plus when it entered it) at every iteration and step, and its load (running, waiting, oldest running sequence, paged KV occupancy) and prefix-cache counts at the top of every iteration, which is also the moment it goes idle, so at rest the numbers are exact. Readers only load atomics.
  • /health is 503 with the same body when the engine thread has exited or the loop overstayed its phase: 30 s between steps, 300 s in a decode step, 1800 s in a prefill. A long prefill is work, not a stall.
  • Counters are process-wide and only grow: engine/ledger.rs statics, fed once per finished sequence from its Usage in finish_or_add_toks_to_seq. An engine rebooted inside the process does not reset them; n > 1 counts each choice.
  • Ledger line: INFO on target hanzo_engine::engine::ledger, message starting serve_api: so bench-serving.py's regex matches it. Rounds are derived, not counted: n - 1 - draft_accepted (a verify round commits accepted + 1, the first token comes from prefill). Undrafted requests log batch with no rounds, which the bench's regex skips, as it skips serial requests.
  • Drafter names come from SpeculativeAttachInfo::name: mtp, dflash, else spec. Pipelines record the attach info (Pipeline::drafter); Hanzo::config carries it.

Request surface: budgets, reasoning, sampling, tools, errors

The rules are pure functions in hanzo-server-core/src/controls.rs (checked in one fixed order by chat_completion::parse_request, which Responses and Messages reuse), errors in error.rs, output shapes in shape.rs.

  • Defaults are hanzo serve flags (defaults.rs, Defaults::validated at startup): a value the request sends always wins; nothing out of range is clamped, it is refused.
  • Temperature: request, else --temperature, else greedy. 0 is greedy.
  • Budget: max_tokens = max_completion_tokens = max_output_tokens (different values are a 400); default --max-tokens (completions --completion-max-tokens); over --max-tokens-cap is a 400; <= 0 is 1. The think-block budget is the request's or --max-thinking-tokens, capped at max_tokens - room (room --thinking-answer-room, else max(1024, 15%)), never 0.
  • Effort reaches the template through one function, chat_template::template_effort: the top level is xhigh on templates that name it (Qwen3.8 raises on anything else), high elsewhere.
  • Errors are OpenAI's envelope everywhere except /v1/messages, which keeps Anthropic's; engine ValidationError is a 400, a declined forced call (FORCED_CALL_DECLINED) a 502, and an error after a stream started is a data: {"error":...} frame before [DONE].
  • Streams: usage is on the stream only with stream_options.include_usage (a choices: [] chunk before [DONE]); the finish chunk carries timings. Responses and Messages read usage from the engine's final chunk themselves, so framing lives only in the chat/completions SSE layer.

Working here

LLM.md is the canonical guide. CLAUDE.md and AGENTS.md are symlinks to it and all three are tracked -- nothing here is gitignored, and the two aliases were real, separately-drifting files until they were folded in.

  1. Record a significant discovery in this file. Never in a new summary file.
  2. Always cargo check before calling an integration complete, and if it does not compile, keep editing. Do not return TODOs.

Code style (extremely important, and the convention here)

Comments. Default to none. Add one only when the why is not obvious from the code: a hidden constraint, an invariant, a surprising edge case, a reference to a spec or an HF source. Never paraphrase the next line, never restate the function name, never narrate steps.

  • Multi-line comments are discouraged in code, and are really only for documentation or where they are the best way to carry the information.
  • One line each, up to ~120 cols. No multi-paragraph /// blocks, no bulleted lists in doc comments, no // === Section === banners.
  • Terse and casual, and never explaining what the code directly below does.
  • ASCII only unless instructed otherwise. No em-dashes, en-dashes, ellipsis characters, smart quotes or box-drawing. Do not use --. ..., " and ' are fine.
  • Do not reference the current task, PR, fix or commit in a comment. That belongs in the PR description, and it rots as the codebase moves.
  • Terse trailing annotations like // already sent above are fine.

Magic values. Hoist durations, sizes and sentinels to named consts at the top of the file. A sentinel that crosses a module boundary -- one place sets Some(0), another checks for it -- must be a pub const, not a literal both sides happen to share.

Function shape. At 6+ arguments, wrap the invariants in a small context struct (DispatchCtx<'a>). Do not add error handling, fallbacks or validation for cases that cannot occur; trust internal code and framework guarantees. No backwards-compatibility shims unless asked.

Pull requests. Never include a "Test plan" section. Conventional commit style: feat(crate):, fix(crate):, docs:. Update docs/src/content/docs/ and the examples alongside a feature or a breaking change.

Test, lint, docs

export HF_TOKEN=<token>      # TESTS_HF_TOKEN for CI parity; some tests need it
cargo test -p hanzo-engine -p hanzo-quant -p hanzo-vision
cargo test --workspace       # crates without tests are skipped
make fmt                     # rustfmt + ruff + clang-format
cargo fmt --all -- --check
cargo clippy --workspace --tests --examples -- -D warnings
cargo doc --workspace        # rendered under target/doc/

CI (.github/workflows/ci.yml) runs exactly: cargo check all targets, cargo test on the core crates, cargo fmt -- --check, cargo clippy -D warnings, cargo doc, and crate-ci/typos. The docs site builds with cd docs && npm run build and deploys through .github/workflows/docs.yml.

Unit tests sit beside their source; integration tests live in tests/; cargo test -p <crate> narrows to one. Python tests need the package built and installed first.

Where each feature lives

Paths are relative to docs/src/content/docs/. The CLI auto-detects model type, so every one of these is hanzo-engine run -m <model>.

Feature Crate Docs
Text inference hanzo-engine, hanzo guides/customize/sampling.md, guides/agents/
Multimodal hanzo-vision explanation/multimodal-pipeline.md, reference/supported-models.md
Diffusion, speech -- reference/supported-models.md
Quantization and ISQ hanzo-quant reference/quantization-types.md, explanation/quantization-tradeoffs.md; scripts/convert_awq_marlin.py
Paged attention hanzo-paged-attn explanation/paged-attention.md, guides/perf/use-paged-attention.md
Adapters, LoRA, X-LoRA -- guides/customize/lora-adapters.md
AnyMoE -- guides/customize/anymoe.md

Also outside the workspace: chat_templates/ (JSON/Jinja), examples/ (Rust, Python, server, notebooks), scripts/.

Two pitfalls that produce wrong output rather than an error

torch.bucketize(right=True) needs Ok(i) => i + 1. Rust's binary_search_by returns Ok(i) at the found position, which is bisect_left. For bisect_right you must insert after equal elements. Err(i) => i is right either way.

Causal Conv1d left padding is effective_kernel_size - stride, not (kernel_size - 1) * dilation -- that is the total padding for the non-causal case. kernel_size=3, stride=2, dilation=1 gives 1, not 2. Check against the HF model's VoxtralRealtimeCausalConv1d or its equivalent.

DeepSeek-V4 + SOTA Inference Playbook (2026-06-28)

Cross-verified research (the V4 config.json and arXiv paper, public reference implementations) + this session's GB10 (sm_121, 273 GB/s LPDDR5X, 128 GB unified) measurements. The binding constraint for batch-1 decode is MEMORY BANDWIDTH (~11 GB read/token); every win = fewer bytes/token or higher BW utilization. Compute-side FP4 tensor-core gains are IRRELEVANT at batch-1 (and broken on sm_121).

The window bug, and what it teaches about this model (2026-09-15)

Found while onboarding the vision variant; it was present in both ports and neither's tests caught it, because they pinned the layer mode and never the window.

  • compress_ratios[il] selects a layer's compressor, not its context. Read as "no window", it gave the ratio-0 layers — and on the safetensors path the ratio-4 Indexed layers too, about half the model — unbounded attention, against a reference that windows every layer (its KV cache for those layers is window_size long).
  • Setting SdpaParams.sliding_window is not enough. Every shipped checkpoint carries attn_sinks on all 43 layers, so run_attention dispatches to the sinks backend before the flash block. head_dim 512 misses every templated sinks kernel, and only the fused CUDA decode branch (q_len 1..=8, no mask) reads window_size. CUDA prefill, Metal and CPU honour the mask alone, so the window must also arrive as a mask or a KV slice.
  • The old path was worse than unwindowed. With flash prefill on by default, CUDA prefill at q_len > 8 produced CausalFlash, which carries no tensor, and reached the unfused sinks path with no mask at all: non-causal attention. combine_with_mask needs a base tensor, so DSA selection was silently disabled on the Indexed layers as well. Any earlier number from that path measured a different model. The GGUF path was unaffected — it has always used CausalMaskConfig::gguf().
  • Forcing a custom mask costs almost nothing here. Flash is unreachable while sinks preempt it, and the eager sinks path already materialises the scores, so the mask adds one L² plane beside 64·L². Passing the window to flash instead only pays off once the sinks backend can serve head_dim 512.

Make it unrepresentable, don't remember it (queued)

The root cause has exactly two instances, both DeepSeek-V4 and both fixed (deepseek4.rs, deepseek4_mtp.rs); the GGUF window bug is a sibling, not the same cause. Counted with the right predicate — sinks: Some(, since grepping sinks also matches sinks: None — only three files in the crate load sinks at all, and the gpt-oss pair is safe because head_dim 64 hits a templated kernel that applies causality itself.

Two instances is not an argument for vigilance. The bug needs three independent facts to line up: sinks present, a head_dim outside every templated kernel's set, and a mask config that may return a tensor-free variant (CausalFlash). No reviewer holds all three at once, which is why it survived tests that pinned the layer mode. The fix is structural: run_attention should refuse a tensor-free mask when it is about to dispatch to a backend that reads only the mask, so the combination fails at startup or at compile time instead of silently producing bidirectional attention. Not attempted during the merge.

V4 architecture — the corrected mental model

  • NOT MLA. V4 replaced MLA with hybrid sparse attention (DSA), per-layer mix of: SWA-128 (raw sliding window) · CSA (4×-compressed KV, lightning-indexer top-k 512/1024) · HCA (128×-compressed KV, dense) · learned per-head attention sinks. The head_dim:512 / kv_heads:1 is a shared-K=V MQA head with partial RoPE-64 — compressed-KV MQA + sparse selection, not MLA up-projection decompression.
  • Weights ship FP8 (E4M3, 128×128 block, UE8M0 µscale) dense + native FP4 MoE experts, bf16 master. The model is W8A8 QAT-trained — it EXPECTS low-precision activations.
  • MoE: 256(Flash)/384(Pro) routed, 6 active + 1 shared; gate = sqrtsoftplus + noaux_tc bias; clamped SwiGLU (limit 10); 3 Hash-MoE bootstrap layers (tid2eid). Residual = mHC (4-stream, 20-iter Sinkhorn doubly-stochastic mixing).

Quant recipe (canonical = what we run; the published IQ2 GGUF, imatrix-built)

IQ2_XXS(ffn_gate/up_exps) · Q2_K(ffn_down_exps) · Q8_0(shared experts, all attn proj, output) · F16(token_embd, router) · F32(norms, sinks) · I32(hash). q2≈81 GiB. imatrix REQUIRED for IQ2_XXS (from real routed-expert activations). Q3=garbage, Q4 decode unstable (argmax collapse from 43-layer error accumulation — the fix is hi-precision matmul accumulation).

The decode-degeneration finding (this session) — WHY hard/long prompts collapse

V4-at-IQ2 sits at a repetition bifurcation: cumulative quant error across 43 layers tips it. The fix is hi-precision matmul accumulation, not higher activation precision. KEY MEASUREMENT: F32 is WORSE than bf16 here — because the model is FP8-trained, so MORE precision moves AWAY from the trained operating point. bit-exact = Q8_0/IQ2 kernels that accumulate in int32 dp4a / hi-precision, NOT go f32. Current best config: bf16 carrier + native dp4a (CUDA_FAST_MMQ / V4 set_fast_mmq(true)) → correct on short/medium (primary-colors, France, math, natural EOS); haiku-class long-hard reasoning still degenerates (the accumulation-precision residual).

Improvement playbook (ranked by batch-1-decode leverage on GB10)

TIER 1 — proven, highest leverage

  1. Per-layer CUDA-graph capture + split-K F16/dp4a decode GEMV. 73→94% of roofline is ~19-25 t/s at 11 GB/token on 273 GB/s (we're at 3.5). Per-LAYER (resolve MoE routing host-side / fixed expert capacity), because a data-dependent expert matmul breaks whole-graph capture. cudarc graph capture. This is the #1 perf item.
  2. dp4a int8 (NOT tensor-core INT8/FP4) for decode GEMV. At N=1, WMMA wastes 15/16 slots (0.25×). FP4 tensor cores only help PREFILL (build sm_121a + CUDA 13 compute_121f). Decode = dp4a / F16 split-K.
  3. Bandwidth-optimal IQ2_XXS dequant→GEMV (wide 128/256-bit loads, in-register dequant, L1::no_allocate on streamed expert weights, evict_last on reused activation, maxrregcount 32-45). The routed IQ2 experts dominate the 11 GB/token. Dequant bit-exact-tested against the GGUF block definitions.
  4. Hi-precision accumulation in ALL quant matmuls (the haiku fix): int32 dp4a accumulate + verify IQ2 mmvq accumulates hi-precision. This is bit-exact correctness for hard prompts.

TIER 2 — correctness + structural

  1. Real V4 sparse attention (DSA: CSA top-k lightning-indexer + HCA + SWA-128 + sinks, FP8 KV) — two caches: latents + indexer keys. The compressor is wired (this session); the indexer top-k selection is the missing piece for >128 context (1M ctx ≈ 9.62 GiB/seq).
  2. Fuse MoE hot path (TopK gate+route, gate/up SwiGLU sharing activation load, RMSNorm+mul, compressor+RoPE+cache-insert).
  3. Unified-memory zero-copy (mmap GGUF GPU-addressable via Grace ATS; cudaMemAdvise ReadMostly/AccessedBy on resident, PreferredLocation(CPU)+AccessedBy(GPU) on cold experts; prefetch predicted experts). Enables SSD/CPU expert streaming.

TIER 3 — conditional

  1. EAGLE-3 speculative (> native MTP) — accept ~2.5 vs MTP ~1.9. MANDATORY: batched verify (stream weights ONCE for all K draft tokens, fixed-shape CUDA graph) else net-negative single-stream (depth-1 MTP measures −21%). The MTP head (3.6 GiB gguf) is being loaded; wire it EAGLE-style, not depth-1.
  2. DeepSeek-OCR — a separate vision+V2 port (SAM/CLIP encoder + DeepEncoder + V2-MoE decoder). Scoped follow-on.

DON'T (negative results, save the effort)

FP4/INT8 tensor cores for decode (batch-1 irrelevant, sm_121-broken) · whole-model CUDA graphs (MoE breaks them → per-layer) · native depth-1 MTP single-stream · DeepGEMM/FlashMLA drop-in (no sm_121 FP4 path mid-2026) · going F32 for "precision" (worse — model is FP8-trained).

Build note

sm_121: build sm_121a + CUDA 13.

Speculative: prompt-lookup (n-gram) proposer -- SHIPPED (engine 1.7.83)

Zero-training, quantization-agnostic draft that reuses 100% of the batched-verify engine (driver/verifier/proposer) with NO verifier change. SpeculativeConfig::PromptLookup { ngram_min, ngram_max, gamma } -> speculative::prompt_lookup::PromptLookupProposer. The draft is the sequence's OWN history: the tail n-gram (longest first, ngram_max..=ngram_min) -> its most-recent earlier occurrence -> the <=gamma tokens that followed. Returns logits: None -> the verifier's exact-match path -> provably lossless (greedy byte-identical on/off). Auto-gates: drafts only at batch <= 4.

  • Attach in pipeline/{normal,gguf}.rs attach_speculative (exempt from PagedAttention, like DSpark) -> draft_proposer. Ref: apoorvumang/prompt-lookup-decoding = HF prompt_lookup_num_tokens.
  • CPU/GGUF now works: the GGUF non-paged spec path was made GENERIC (was Deepseek-only) so any normal-cache GGUF (qwen2/qwen3) drives NormalSpeculativeCacheAccess. CLI --prompt-lookup-ngram <N> (reuses --gamma); server default-OFF (opt-in). Gate test: hanzo-cli/tests/prompt_lookup_lossless.rs (CPU).
  • MEASURED (zen-eco-4b Q4 GGUF, CPU, grounded code-repeat prompt): 85% acceptance (mean_accepted 6.83/8), byte-identical. ~6 target forwards emitted ~47 tokens vs 48 forwards without = ~8x fewer target forwards. Wall-clock ~break-even here (0.85-1.12x across runs) because CPU decode of this small model is COMPUTE-bound: a width-(1+gamma) verify forward costs ~(1+gamma)x a width-1 one, so spec-decode can only win on MEMORY-bound decode (GPU, or large models streaming weights once per forward -> wide verify ~free -> throughput ~= 1+accepted). Same regime law as the EAGLE/MTP note ("stream weights ONCE for all K").

zen6 MTP head (models/quantized_zen6_mtp.rs)

  • The head is the MTP GGUF published beside the model (mtp-Qwen3.8-Flash-Next-{shared-,}Q8_0.gguf, arch qwen4exp, read as zen6, block_count 49, nextn_predict_layers 1). One attention+MoE block at blk.48, behind the same 4-stream hyper-connections as the trunk. Entry: eh_proj [enorm(embed(x_{t+1})) | hnorm_s(h_t,s)] per stream, where h is the trunk's wide streams before the output mixer. Exit: its own mixer (nextn.hc_head_*), then lm_head.
  • shared files lack token_embd/output and borrow the target's. Measured: the self-contained file's token_embd is byte-identical to UD-IQ4_XS's; its output is Q8_0 where the target's is Q6_K.
  • The head's attention keeps a KV over the whole sequence. Every target forward writes its pairs' K/V (the hook in ModelWeights::forward), keyed by sequence id from note_forward_sequences; a draft step truncates at its own position, which drops rejected pairs. The Qwen3.5 and DeepSeek heads here keep only their chain (reset every round): at depth 1 that is one key, so attention returns V and the head drafts context-free. Expect lower acceptance from those.
  • --prompt-lookup-ngram N --gamma G beside --mtp-model continues each head draft by lookup (speculative::Extended): a needle of exactly N tokens ending in the drafted token adds up to G tokens. Greedy batches only.

CUDA flash prefill: the missing-.contiguous() bug -- FOUND + FIXED (573d12614)

  • SYMPTOM: CUDA GGUF prefill was BOTH slow (eager O(n^2) collapse: pp128 2158 -> pp512 2044 -> pp2048 1118, where a working flash path is flat) AND, when routed to flash, produced GARBAGE logits (byte-identical 钊，，， from token 1 -> corrupted KV cache -> whole generation garbled).
  • ROOT CAUSE (one bug, two symptoms): the can_use_flash path in Sdpa::run_attention transposed q/k/v (b,H,s,d)->(b,s,H,d) WITHOUT .contiguous(). flash_attn_v2 hands the tensor straight to the CUDA kernel with no internal contiguous, so for seq>1 (prefill) the non-contiguous strides are read wrong -> garbage. Decode (seq==1) is trivially contiguous so it worked -- which HID the bug and made it look prefill-specific. Because flash garbled, GGUF prefill had been left on the eager naive_sdpa fallback (the O(n^2) collapse). ONE .contiguous() fixes both.
  • HOW IT WAS FOUND (methodology, not guessing): (1) an ISOLATION TEST (flash_correctness in attention/ mod.rs) proved flash_attn == naive_sdpa to maxabs 0.0156 on controlled bf16 GQA-causal inputs WITH .contiguous() -> the crate/kernel is CORRECT (killed the "Blackwell bf16 kernel bug" theory). (2) a DBG_ATTN trace showed the real prefill call is mask=CausalFlash q_dims=[1,32,512,128] -> the can_use_flash path (NOT the Custom-mask block I'd been chasing). (3) that path lacked the .contiguous() the passing isolation test had. Fix = add it. Earlier guesses (Custom-mask route, GQA repeat_kv, window, bf16 dispatch) were all on the WRONG path -- the trace corrected course.
  • RESULT (8B CUDA, spark GB10, main + fix): prefill coherent + FLAT -- pp512 2653 (pure) / 2349 (with decode); pp2048 ~2586 flat vs the old eager 1118 (the big long-context win). Decode 35.8. The pp2048 collapse is gone. Fixes flash prefill for ALL CUDA models (any that hit the CausalFlash path), not just GGUF. What prefill time remains is kernel efficiency, not an algorithmic O(n^2) defect.
  • Flash dispatch: using_flash_attn() (CUDA) is now default ON with a FLASH_PREFILL=0 opt-out; flash is used iff that predicate AND applicable (device/dtype/shape/causal). ROCm flash is always-on when applicable.

8B repetition-collapse ROOT-CAUSED: flash PREFILL precision (a regression from the .contiguous fix)

  • SYMPTOM: Qwen3-8B-Q4K_M on CUDA (and Vulkan) collapses to 钊，，， repetition during LONG thinking-mode generation (plain prompt -> <think> block). Short / /no_think output is coherent. Metal is fully correct (same byte-identical GGUF -> "Paris"). 0.6B/4B fine on all backends.
  • LOCALIZED by A/B (seed 42, plain prompt): FORCE_EAGER (prefill+decode eager) = COHERENT; DECODE_EAGER (prefill flash, decode eager) = STILL COLLAPSES. So the collapse originates in flash PREFILL, not decode. NOT the file (identical sha256), NOT the quant kernel (DEQUANTIZE_ALL still collapses).
  • ROOT CAUSE: the CUDA flash-attn prefill is ~1.5% numerically off vs eager naive_sdpa (isolation test flash_matches_naive maxabs=0.0156 -- larger than a correct flash's ~0.4% bf16-reorder). That error, accumulated across 36 layers, tips 8B's sensitive thinking-mode trajectory into the repetition bifurcation. Eager (f32 softmax) is exact and stays stable; Metal's attention is a different, stable impl.
  • THE TRADE-OFF THIS EXPOSES: the .contiguous() flash-prefill fix (573d12614) gave prefill 0.37x->0.86x flat BUT introduced this collapse (pre-fix, GGUF prefill used eager = correct-but-slow). So main now ships fast-prefill + 8B-long-gen-collapse on CUDA/Vulkan. Options: (a) revert to eager (correct, slow prefill); (b) fix flash's softmax to f32-accumulate precision (deep, in hanzo-flash-attn -- the real fix, keeps perf); (c) the kernel-DSL migration (one correct attention impl across backends -- the structural cure). DECODE_EAGER is a NEGATIVE (does not fix) -- do not ship it.
  • NEXT: diagnose whether hanzo-flash-attn accumulates the softmax in f16 vs f32 (the likely 1.5% source) and force f32. If the flash kernel is correct, the 1.5% is the is_causal alignment -- verify the causal triangular direction matches naive. Until fixed, 8B-long-thinking on CUDA/Vulkan is a known-issue (usable for short/direct gen; Metal fully correct).

RESOLVED: flash prefill default ON (dense) + GGUF gated (spark GB10, hanzo-flash-attn 0.11.35)

  • using_flash_attn() flipped default ON (FLASH_PREFILL=0 opts out) after bumping hanzo-flash-attn to 0.11.35 (routes bf16 softmax-P through the f16 kernel's 10 mantissa bits vs bf16's 7).
  • DENSE bf16 8B on CUDA: flash prefill is BYTE-EXACT vs eager (greedy seed-0, first-prefill non-varlen path) and 1.34x@pp512 / 2.46x@pp2048 faster (flat vs length: eager 2110/1166, flash 2825/2868 T/s). Decode 14.8->15.0 (paged kernel, unchanged). d_flash == d_bf16raw == d_eager, so the P-cast is NOT what tips dense here.
  • GGUF/quantized garble under flash and are GATED to eager via CausalMaskConfig::gguf() (all quantized_* models). Verified: Qwen3-8B-Q4K flash emits corrupted text ("-devel...corrupted") that is P-precision-INDEPENDENT (q_flash == q_bf16raw), so it is NOT the 1.5% P-cast issue above -- it is structural: GGUF forward_attn has no FlashParams plumbing (passes None), so flash runs without cumulative seqlens. Q4K + MoE flash == eager (coherent "Paris") after the gate. This is the same config as the "8B-Q4K thinking-mode collapse" above.
  • The paged prefix-cache/chunked GATHER path (varlen, q_len<kv_len) is coherent under flash but its greedy output differs from a fresh (non-cached) run -- and EAGER gather differs from fresh too, so this is a pre-existing prefix-cache non-transparency (reduction-order/fp between gather vs non-gather kernels), NOT a flash-specific garble. Single-shot generation never hits the gather path (needs a prefix-cache hit), so the default-ON win is safe there. A transparent gather (build-2) remains a separate follow-up.

CORRECTION: the sampling change is NOT the 8B collapse fix (drill continued)

  • Tested the sane-sampling fix (rep_penalty 1.1, top_p 0.9, temp 0.7) on 8B seed42 flash prefill: STILL COLLAPSES -- 针needle needle needle... (different token than 钊，，，, but still a repetition attractor). So the collapse is NOT a sampling artifact -- rep-penalty does not rescue it.
  • CONCLUSION (all hypotheses eliminated by A/B): the collapse is hanzo-flash-attn's PREFILL being subtly off. Its softmax is correctly f32, yet its output (~1.5% vs eager) tips 8B-Q4K bf16 into a repetition attractor -- while EAGER and Metal's attention both keep 8B stable. So hanzo-flash-attn has a subtle numeric error, NOT just "expected reorder".
  • The sampling change is kept as a genuine DEFAULT improvement (temp 0.1/top_p 0.1/no-penalty was bad), but it is NOT the collapse fix. Do not present it as such.
  • REAL FIX PATHS: (a) revert flash prefill -> eager (correct, prefill 0.37x -- correctness-first); (b) fix hanzo-flash-attn's CUDA kernel by numeric-diffing its prefill output against the eager path's to find the subtle error (deep, keeps the 0.86x perf); (c) the kernel-DSL migration (one correct attention impl -- structural cure). Production (server/API, client-set sampling, short gens) is largely unaffected; the collapse is a bare-CLI long-thinking-gen issue on CUDA/Vulkan.

FINAL VERDICT (numeric diff done): 8B collapse = bf16 BIFURCATION SENSITIVITY, not a fixable flash bug

  • The flash-vs-eager seq sweep (attention::flash_precision_probe, CUDA, bf16 GQA causal 8B-shape) settles it: seq=8 max_abs=0.0156 | seq=32 0.0156 | seq=128 0.0156 | seq=512 0.0176 The error is FLAT with seq at ~0.0156 = ONE bf16 ULP for these magnitudes. So it is NOT a systematic algorithmic bug (which a short-seq test would also show, but a real bug would be a LARGER fixed offset) and NOT reduction-reorder (which would SCALE with seq). It is the bf16 OUTPUT-QUANTIZATION floor: flash and eager compute slightly-different-but-both-VALID f32 attention outputs that round to adjacent bf16 values. flash's softmax is f32-correct (verified: MaxOp/SumOp/ElementAccum=float).
  • CONCLUSION: there is NO clean flash bug to fix. 8B-Q4K at bf16 sits on a repetition bifurcation; flash's valid rounding tips it, eager's valid rounding does not (Metal is stable because it uses its OWN eager/ steel attention, NOT the CUDA flash crate). Better sampling did not rescue it (rep-penalty still looped).
  • SHIPPED RESOLUTION (pragmatic, CTO call): KEEP flash prefill on main -- the 0.86x-flat prefill win is real and broad; correct for the vast majority of inputs; production (server/API, client sampling) is unaffected; the collapse is a NARROW bare-CLI long-thinking-gen edge on the one sensitive 8B-Q4K model. The CLI sampling defaults were also improved (temp 0.7/top_p 0.9/rep-pen 1.1 -- objectively better than the old temp 0.1/top_p 0.1/no-penalty) though they do not fix the bifurcation. flash_vs_eager_seq_sweep stays as the regression gate.
  • THE PERMANENT CURE is the kernel-DSL migration (hanzo-kernel, proven CPU+Vulkan+Metal this session): ONE correct attention implementation lowered to every backend eliminates the flash-vs-eager-vs-Metal numeric fork by construction -- the whole "same op, N impls, N bf16 roundings, N bifurcation outcomes" class.

Pixal3D: native image-to-3D (TRELLIS-image-large, MIT) -- diffusion_models/pixal3d/

  • Target = microsoft/TRELLIS-image-large (the mature 1.2B model). TRELLIS.2-4B (DINOv3, ConvNeXt VAE, RoPE, shape/texture cascade) is the successor and a 3x heavier follow-up, NOT this port.
  • Pipeline stages: image -> DINOv2 tokens -> sparse-structure rectified-flow -> Conv3d decode to a 64^3 occupancy grid -> [SLAT sparse flow -> FlexiCubes mesh] -> GLB. The DENSE half (through occupancy) is fully ported; the SLAT/FlexiCubes half is the frontier (needs sparse-conv infra + FlexiCubes).
  • ALL dense stages are BIT-EXACT vs the torch reference oracle (cos=1.00000000): dinov2 (max|d|=8.2e-5), ss_flow (mse=3.0e-12), ss_decoder (mse=2.3e-9). Oracle = TRELLIS's own pure-torch code run CPU-only with ATTN_BACKEND=sdpa and the sparse/render subpackages stubbed (sys.modules['trellis.{pipelines,renderers,representations}']=ModuleType(...)), so no spconv/nvdiffrast/flash needed to validate the dense models. Scripts in scratchpad/trellis_oracle.
  • Key facts pinned from the reference source (do not re-derive):
    • DINOv2 = torch-hub dinov2_vitl14_reg (NOT the HF-naming variant): fused qkv, LayerScale ls1/ls2, 4 registers inserted AFTER cls (no pos-emb), 518->1374 tokens, exact-erf GELU; TRELLIS uses x_prenorm (before the final norm) then a non-affine F.layer_norm (eps 1e-5). Block norms eps 1e-6.
    • ModulatedTransformerCrossBlock (SS + SLAT share it): norm1/norm3 non-affine, norm2 affine; self-attn is adaLN-gated with per-head qk-RMSNorm (= L2-normalize over head_dim * gamma[H,D] * sqrt(D)); cross-attn to the DINOv2 tokens is un-gated/un-modulated; FFN is tanh-GELU, adaLN-gated. adaLN = SiLU->Linear(C,6C).
    • Sampler passes t*1000 to the model; timestep_embedding is cos-then-sin; SS pos_emb is a stored buffer.
    • FlowEuler: t_seq=linspace(1,0,steps+1) reparam by rescale_t=3; step x-=（t-t_prev)*v; CFG only inside cfg_interval [0.5,1] as (1+s)cond - suncond, s=5; neg_cond = zeros_like(cond).
  • hanzo-ml has NO Conv3d: the decoder's only kernel (3x3x3, pad1, stride1) is decomposed into 3 depth-sliced conv2d summed with depth alignment (ss_decoder.rs::Conv3d). hanzo-ml's tuple Shape/Dims cap at rank 6 -> use Vec/&[usize] for the 8-D pixel_shuffle_3d reshape/permute.
  • Weights: PIXAL3D_MODEL dir holds the TRELLIS ckpts (ckpts/*.safetensors) + a converted dinov2_vitl14_reg.safetensors (torch-hub .pth has no safetensors; convert once). Runs Device::Cpu today.
  • GLB export = hand-rolled binary glTF 2.0 (glb.rs), validated against the third-party gltf loader; wired into ThreeDFormat::Glb (was a PLY fallback). The /v1/3d async endpoint already calls pixal3d_generate.
  • SLAT + FlexiCubes half is now DONE (the fine textured mesh), all bit-exact vs the reference:
    • sparse.rs = the sparse-tensor infra (SparseTensor coords[N,3]+feats[N,C], B=1): submanifold Conv3d (SubMConv3d, gather active neighbours c+(kz-1,ky-1,kx-1)), SparseLinear, downsample/upsample/ subdivide, SparseGroupNorm32, AbsolutePositionEmbedder, window partition. GOTCHA: spconv's CPU build is ~2% WRONG on sparse SubMConv (single-tap dense is exact, multi-tap sparse gather-gemm is buggy) -- the correct oracle is dense conv3d @ active sites (== GPU spconv). GOTCHA: SparseDownsample uses torch.scatter_reduce(mean, include_self=True) so the divisor is (group size + 1), not the size.
    • slat_flow.rs (SLatFlowModel) reuses the SS-flow ModulatedCrossBlock verbatim (full attention over the voxel set == dense sdpa at B=1) + SparseResBlock3d pack/unpack (downsample 64->32, 24 blocks, upsample 32->64 with U-Net skips). NO stored pos_emb (on-the-fly AbsolutePositionEmbedder from coords). Parity cos=1.0000, max|d|=5.2e-5.
    • slat_decoder.rs (SLatMeshDecoder): 12-block sparse SWIN transformer (windowed self-attn, alternating shift 4, non-affine norms), 2 SparseSubdivideBlock3d (GroupNorm32+subdivide+conv residual) 64->128->256, out_layer -> 101-ch FlexiCubes layout (sdf 8 + deform 24 + weights 21 + color 48). Parity cos=1.0000.
    • flexicubes.rs (+ flexicubes_tables.rs) = nvdiffrec FlexiCubes dual-MC, inference path. Ported SPARSELY: the SDF grid is 1 everywhere except active-voxel corners, so only cubes touching an sdf<0 vertex can be surface -- no dense 256^3 arrays. sparse_cube2verts corner-mean, DMC case + check_table ambiguity inversion, beta-weighted dual verts from alpha-weighted edge zero-crossings, 6-ch color interp, quad triangulation split by gamma with sdf winding. Cubes processed in reg_c (z,y,x) order so the 4-cube quad winding matches. Parity vs golden mesh: verts 32981==32981, faces 63978==63978, Chamfer=0.000000.
    • pipeline.rs generate_textured(): coarse occupancy -> active coords (argwhere>0, res 64) -> SLAT flow denoise (from noise, CFG) -> denormalize (*std+mean, consts SLAT_MEAN/STD) -> mesh decoder -> FlexiCubes. /v1/3d exposes it via the texture request flag; GLB carries per-vertex RGB as COLOR_0 (mesh_to_glb_colored).
  • Oracle method for the SLAT half: real TRELLIS + real CPU spconv, with a flash_attn shim (sdpa CPU) so the real code runs, trellis.{pipelines,renderers,gaussian,octree,radiance_field} stubbed, kaolin.check_tensor stubbed, FlexiCubes forced device=cpu, and SparseConv3d monkeypatched to the correct submanifold gather (CPU spconv is buggy). Scripts in scratchpad/oracle (gen_oracle.py + flash_attn.py).

LLaDA: diffusion LLM (dLLM) native port -- models/llada.rs

  • Model: GSAI-ML/LLaDA-8B-Instruct (MIT, ungated, most-downloaded dLLM). Same code loads LLaDA-1.5 + LLaDA-8B-Base (identical arch). Rejected: DiffuCoder (apple-amlr research-only); Dream-7B (Apache, qwen2- arch) is the fallback; LLaDA-MoE-7B-A1B (Apache) is future.
  • Arch = Llama with 2 differences: BIDIRECTIONAL attention (no causal mask) + iterative masked-diffusion generation (no KV cache, whole sequence re-forwarded each step). OLMo config names (d_model 4096, n_heads 32, n_kv_heads 32 = MHA, n_layers 32, mlp_hidden_size 12288 SPLIT SwiGLU ff_proj gate + up_proj, ff_out down, rms_norm, rope_theta 500000 NeoX, untied lm_head = transformer.ff_out, vocab 126464, mask 126336). Reuses RmsNorm/Sdpa/RotaryEmbedding/ReplicatedLayer verbatim; the ONLY structural change vs llama.rs is the mask. Bidirectional = AttentionMask::None + FlashParams::empty(false) (pitfall 6). Weight prefix is model.transformer.{wte,blocks.N.*,ln_f,ff_out}.
  • Sampler (official LLaDA generate.py, ported): semi-autoregressive blocks (block_length), low-confidence remasking (confidence = max softmax prob = softmax at argmax), linear noise schedule (num_transfer_tokens), greedy (temp 0). Per step: full forward -> argmax x0 + confidence -> unmask the top-k highest-confidence masked positions within the current block. Selection on CPU (tiny), forward on GPU (dominant). Deterministic.
  • PARITY (GB10, bf16, vs transformers-4.46 oracle -- tf5.x breaks LLaDA remote code via all_tied_weights_keys): gate1 single-forward logit cosine > 0.99; gate2 generation BYTE-IDENTICAL on deterministic prompts (France ->"Paris" [65926,eot,eos...] 32/32; 17*23->"391" [18,24,16,...] 32/32); haiku 19/32 = legitimate bf16 confidence-tie nondeterminism (both valid). Coherence: "Paris", coherent haiku.
  • THE DISTRIBUTED THESIS, measured (gen=128, block=32, GB10 bf16): tok/s scales LINEARLY with tokens-per- traversal (steps 128->1.0 tok/traversal->6.9 T/s; 64->2.0->13.6; 32->4.0->27.7). Each denoise step = ONE model.forward = ONE ring traversal (pp_head_forward). So a 2-node ring PP dLLM amortizes the traversal latency over MANY tokens per hop -- inverting AR's 1-token-per-traversal economics. PP wiring is a thin adapter: PipelineParallelModel = {pp_embed=wte, pp_run_local=local block range, pp_norm_head=ln_f+ff_out}; loop pp_head_forward N times = N traversals. (Trait in pipeline_parallel.rs; single-node sweep proves the economics; 2-node execution + /v1 serving pipeline are the remaining wiring.)
  • No-KV-cost caveat: every step re-forwards the whole [prompt+gen] sequence, so per-step cost grows with sequence length (Fast-dLLM-style block KV cache is the known follow-up). gated tests: llada_parity, llada_smoke (LLADA_WEIGHTS + LLADA_ORACLE / LLADA_GEN/STEPS/BLOCK).

ACE-Step music: /v1/audio/music (this session)

  • Made the composable AceStepPipeline (diffusion_models/ace_step/) servable as a new SpeechLoaderType. UMT5-base text encoder + DiT flow-match (APG guidance) + DCAE + HiFi-GAN vocoder -> stereo 44.1kHz.
  • Registration mirrors TTS exactly: SpeechLoaderType::AceStep (FromStr ace_step|ace-step|acestep, auto-detect on the DiT config's _class_name == "ACEStepTransformer2DModel"), SpeechGenerationConfig:: AceStep { frames, steps, guidance_scale } (default ~10s = 10*44100/4096 frames, 27 steps, guidance 15). Reuses ModelCategory::Speech + Response::Speech + RequestMessage::SpeechGeneration -> ForwardInputsResult ::Speech; forward_inputs tokenizes with the UMT5 tokenizer (umt5-base/tokenizer.json), runs generate(), and interleaves the (1,C,S) waveform into PCM (channels=2, rate=44100).
  • Loader (pipeline/speech.rs) fetches the 4 real ACE-Step sub-checkpoints from ACE-Step/ACE-Step-v1-3.5B: umt5-base/model.safetensors (keep all), ace_step_transformer/diffusion_pytorch_model.safetensors (drop the unused lyric*/projectors/.add_/.to_add_out heads), music_dcae_f8c8/... (keep decoder.*), music_vocoder/... (keep all). All loaded f32 (vocoder/DCAE fidelity). Key filters byte-match the proven ace_step_generate_e2e test fixtures; the DCAE f8c8 gives 4096 samples/latent-frame (44100/4096 ~10.77 fps).
  • Server: POST /v1/audio/music (MusicGenerationRequest {model,input,response_format}) in hanzo-server-core/src/music_generation.rs reuses speech_generation's parse + response plane (wav/pcm). Merged to main (5c8b8968b). Per-request duration/steps/guidance = config-driven (load-time), same as TTS; threading them per-request would need RequestMessage::SpeechGeneration to carry gen params (follow-up).

Customer surface: where a Hanzo SaaS customer generates audio/video/3D

  • Engine (this repo) serves the OpenAI-adjacent /v1 modality plane: /v1/chat/completions (text+dLLM), /v1/images/generations, /v1/audio/speech (TTS), /v1/audio/music (ACE-Step), /v1/audio/transcriptions (ASR), /v1/3d (TRELLIS/Pixal3D, texture flag), /v1/videos (WAN async), /v1/animate (MuseTalk dub). Deployed via platform.hanzo.ai on DOKS behind the gateway (api.hanzo.ai).
  • studio.hanzo.ai (hanzoai/studio, ComfyUI fork; deployed by hanzo.yml onto do-sfo3-hanzo-k8s / studio Service CR) is the PRIMARY authoring surface: the hanzo_engine node pack (custom_nodes/hanzo_engine) maps one node == one /v1 endpoint over a shared client; the flagship workflow hanzo-full-generative- pipeline.json chains Chat -> ImageGen -> (textured ImageTo3D + WAN video) + TTS + Music.
  • hanzo.app (app builder + /games) is the consumer/preview surface for the produced assets.
  • PRODUCTIZATION GAP (hosted multi-tenant), named precisely from the studio code: (1) Engine URL is not per-tenant: the node pack reads process-global HANZO_ENGINE_URL (default localhost:1234); it is set NOWHERE in the studio repo, and the per-org engine/worker routing (middleware/engine_selector.py + compute_config.py + prompt_router.py, which route the graph to local vs a BYO-GPU worker) never derives or sets the node pack's /v1 base URL per org/worker. (2) Identity is not propagated to the engine: iam_auth_middleware.py authenticates the Studio surface (Hanzo IAM JWT, org from the owner claim) and worker_client signs Studio->worker calls, but custom_nodes/hanzo_engine/client.py sends the /v1 calls with NO Authorization header -> a gateway- fronted engine (api.hanzo.ai injects identity, scopes by owner) gets an anonymous/unscoped call. (3) Hosted GPU + model availability: /v1/3d,/videos,/audio/music,/animate are GPU-bound and need those specific models loaded; the compute layer schedules the ComfyUI graph on GPU workers but there is no per-org binding of "engine deployment X (with ACE-Step/WAN/TRELLIS/MuseTalk loaded) serves this org". Self-hosted single-engine works TODAY (set HANZO_ENGINE_URL, run hanzo-engine serve); the gap is the 3-part wiring to make it hosted SaaS: per-org engine-URL resolution from engine_selector, IAM-token forwarding onto the node /v1 calls, and per-tenant engine/model provisioning.

FP8 -> GGUF converter for GLM-5.2 (glm-dsa) -- scripts/convert_fp8_to_gguf.py

  • WHY: unblock loading GLM-5.2 in-engine. The loader is models/quantized_deepseek2.rs, which accepts general.architecture in {deepseek2, glm-dsa}; GLM-5.2 = glm-dsa (split-MLA + 256-expert MoE).
  • DISK-SAFE: stream one ~5 GB FP8 shard, requantize, delete it, next. But GGUF's whole tensor directory precedes its data blob, so it can't append per shard. So: (1) PLAN the full directory + metadata from config.json, write header/KV/tensor-info, preallocate; (2) stream shards, pwrite each tensor (or one expert's slice of a rank-3 bank) into its fixed offset. Resumable via a .progress.json sidecar keyed by a plan signature (offsets shift if any --*-type/config changes).
  • THE glm-dsa GGUF CONTRACT (matched byte-for-byte to the reader + utils/gguf_metadata.rs):
    • routed experts -> rank-3 blk.{i}.ffn_{gate,up,down}_exps.weight, ne=[hidden, moe_inter, n_expert], Q4_K.
    • router ffn_gate_inp.weight + no-aux exp_probs_b.bias -> F32 (reader dequantizes them).
    • split-MLA: attention.key_length=qk_nope, key_length_mla=q_head_dim, value_length_mla=v_head_dim, rope.dimension_count=qk_rope. q_lora path uses attn_q_a/attn_q_a_norm/attn_q_b; kv_b shipped combined.
    • expert_gating_func=2 (sigmoid), expert_group_count=1 (noaux), leading_dense_block_count=first_k_dense.
    • MTP/nextn block at index n_layers, emitted Q8_0 -- an int4 draft head measures ~0% acceptance, so speculation never starts. nextn_predict_layers records it; the text loader drops it from block_count (n_layers = block_count - nextn).
    • VALUE-TYPE gotchas from hanzo-ml gguf Value: to_f32 accepts ONLY F32 (rms_eps/rope.freq_base/ expert_weights_scale must be FLOAT32, not F64); to_bool ONLY Bool; to_u32/to_u64 upcast U8/U16/U32. context_length is .unwrap()ed by the device mapper -- must be present.
  • QUANT: gguf-py 0.19 quantizes Q8_0/F32 but NOT the K-quants (it does dequantize Q4_K). So Q4_K is a faithful quantize_row_q4_K_ref (make_qkx2 + K_SCALE packing, numpy-vectorized over super-blocks), proven by round-tripping through gguf-py's own Q4_K dequantizer (--selftest, relerr ~0.075).
  • VALIDATED here without the 756 GB checkpoint: --selftest builds a tiny synthetic glm-dsa GGUF end-to-end and asserts every reader-required key/type/tensor/shape via gguf-py GGUFReader; an on-disk FP8(e4m3 128-block)+BF16 safetensors fixture exercised the --indir streaming path and resume. NEEDS a real zai-org/GLM-5.2-FP8 checkpoint for full-scale conversion + an in-engine load smoke test (and the tokenizer is loaded separately -- GGUF vocab embedding is not done here).

Agent session placement in the native router

hanzo-router/src/scheduler.rs extends the existing replica pool for the HTTP front: scoped session pins, successful-response prefix locality hints, configured worker roles/capacity/throughput weights, atomic admission, and observed TTFT EWMA. The legacy ReplicaSet::pick ring API remains available; the HTTP front uses pick_agent. New sessions require a slot; pins tolerate 2× capacity before reassignment. A session stays on the failover worker after recovery. Pins expire at 30 minutes, prefix hints at five, each bounded to 10,000 rows per model. X-Session-Id / X-Agent-Role / X-Target-Replica or JSON metadata session_id / agent_role / target_replica carry hints; keys are SHA-256 of the principal/model namespace and structured session/prefix input, never stored raw prompts. The listener defaults to loopback and trusts identity from the private cloud gateway, not public clients. Its replica admin surface is private too. Cloud's existing metered relay preserves JSON metadata but not arbitrary headers; cloud owns the composition in native/engine-router and deploy/engine-pool.yaml. No KV migration or engine-reported cache/queue telemetry is claimed: completed 2xx response prefixes are locality hints. Streams release leases on EOF/error/drop; only pre-connect errors retry, never timeouts or partially delivered streams.

High-Availability 2-Node GPU Cluster & Streaming Keep-Alive

1. Dual-Node GPU Topology & Failover

  • Spark: Primary engine (NVIDIA Blackwell GB10 NVFP4 on :30000, 128GB VRAM). Designated for Main Thread (TTFT 133ms, 1M context capacity).
  • Evo: Secondary engine (AMD Strix Halo Ryzen AI Max+ 395 Q6_K on :8080, 128GB unified memory). Designated for Subagents (TTFT 618ms).
  • Both nodes run hanzo-router on port :1235 with bidirectional awareness:
    • If Spark fails: router automatically evicts Spark and routes Main Thread to Evo.
    • If Evo fails: router automatically spills Subagents to Spark.
    • In Kubernetes, hanzo-router endpoints include both nodes on :1235 behind hanzo-ingress (Traefik). Addresses live in the private record, not here.

2. SSE Keep-Alive & Mid-Response Connection Protection

  • Problem: During long reasoning or deep thinking cycles (e.g. 10–18 minutes), upstream generation pauses between tokens or during prefill caused NAT firewalls and client watchdogs (CLAUDE_STREAM_IDLE_TIMEOUT_MS, default 300s) to abort connections with "API Error: Connection lost mid-response".
  • Solution in hanzo-router/src/proxy.rs:
    • LeasedStream maintains a 15-second timer. If no data arrives from upstream for 15s during an active SSE stream, it emits an SSE comment (: keep-alive\n\n).
    • Standard SSE clients silently ignore comment lines starting with :, while the socket receives live TCP packets that reset idle watchdog timers and keep stateful NAT tables alive.
    • response_with_progress covers the other half: an engine that sends no HTTP headers until prefill completes gets the same pings, and an upstream error arriving after the first ping is delivered as an SSE event: error. Verified through Cloudflare: a 188K-token request held the connection for 244s of pings, then streamed.
    • Upstream reqwest::Client: 30s connect timeout, 15s TCP keepalive, TCP nodelay, 4s pool idle timeout. The replicas' HTTP servers (uvicorn, cpp-httplib) drop idle keep-alive connections at 5s, which is also the probe period, so a longer pool timeout hands the client a socket the engine is closing; those failures evicted healthy engines every few cycles.
    • Client workstations (ra and dbc) configure API_TIMEOUT_MS=3600000 and CLAUDE_STREAM_IDLE_TIMEOUT_MS=3600000.

2b. Context-capacity admission, liveness and vision routing (2026-09-15)

  • Admission by tokenizer count. Each replica declares max_context in router-pool.yaml; a request is admitted only where prompt + requested output fits. Counts come from each replica's own /v1/messages/count_tokens (pool members serve different models and templates, so counts are not interchangeable), falling back to a byte estimate when a tokenizer cannot answer. Session pins, cached prefixes, roles and explicit targets all yield to capacity. When nothing fits, the client gets invalid_request_error "prompt is too long", which Claude Code answers by compacting.
  • Liveness is /v1/models, not /health. spark's /health runs a one-token generation through the scheduler, so behind a long prefill it answers in seconds; a 2s probe evicted spark, the only replica that could hold a long conversation, and compaction got "no available replica". Probes now ask /v1/models with a 10s ceiling, once per upstream per cycle (not once per model pool), concurrently, and ask twice before declaring a server down.
  • Vision. A replica declares vision: false when its engine has no image encoder (a GGUF replica without a vision projector answers 500). Image-bearing requests — Anthropic image blocks including inside tool results, OpenAI image_url, Responses input_image — skip those replicas and are not token-counted there.
  • tool_reference blocks. ToolSearch results carry them; local engines 500 on them and poison every later turn, so the router rewrites them to text.
  • [1m] model suffix. Claude Code sizes its context window from the model name: with a behavesAs picker row it uses the mapped model's window and ignores CLAUDE_CODE_MAX_CONTEXT_TOKENS, so only a [1m] suffix yields 1M. The client strips the suffix before the request and sends the context-1m beta header instead; the router strips it too so the pool still resolves.
  • Drain needs systemd's permission. The router waits up to 30 minutes for in-flight work on SIGTERM, so both boxes carry ~/.config/systemd/user/hanzo-router.service.d/drain.conf with TimeoutStopSec=1800. Without it systemd SIGKILLs at 90s and the drain is decorative.
  • Restarts are not free. Restarting a router kills in-flight requests, including ones still in the token-counting phase that hold no lease, so an "inflight == 0" drain check is not sufficient. Restarting an engine empties its prefix cache: appended turns normally reuse 166-177K cached tokens and prefill only 17-1,425 new ones, but after a restart or an aborted request the retry reused just 16,384 of a 154K prefix (a hybrid model's recurrent state is kept only at tracked points).

3. ModelOpt NVFP4 in hanzo-quant -- native W4A16, weights stay packed (2026-09-15)

  • hanzo-quant/kernels/nvfp4/nvfp4_gemm.cu expands E2M1 in registers off the same LUT and __byte_perm trick MXFP4 uses. Only the scale differs -- FP8 E4M3 per 16 against E8M0 per 32 -- so one uint4 load carries two blocks and two scales, and e4m3_to_float folds the checkpoint's FP32 weight_scale_2 into each block scale. M <= 4 takes the vecmat decode path, wider shapes a register-tiled one; K must be a multiple of 32. NVFP4Layer keeps weight/weight_scale packed on CUDA (Weight::Packed) and still dequantizes at load everywhere else.
  • Gate: nvfp4_matmul_matches_dequantized on a GB10 compares against dequantize-then-matmul at N/K of 64/128, 512/5120 and 256/17408 for M in {1, 2, 4, 17, 64, 512}.
  • Why it is a correctness fix and not a tuning one: dequantize-at-load turned a 15 GB checkpoint into a 104 GB device-map plan on a 121 GB box and died at CUDA_ERROR_OUT_OF_MEMORY. Packed, the same model loads in ~31 GB.
  • A ModelOpt checkpoint is mixed, per module. RadixArk/Qwen3.8-27B-NVFP4-BF16-LMHead ships 193 NVFP4 modules (the MLPs: weight U8 [N, K/2], weight_scale F8E4M3 [N, K/16], weight_scale_2 F32) and 208 FP8 ones (linear_attn.in_proj_qkv, in_proj_z, out_proj: weight F8E4M3 [N, K], per-tensor F32 weight_scale). load_modelopt_linear routes both, so anything that reads a weight with a raw vb.get instead of the loader gets fp8 bytes and a matmul dtype mismatch -- which is exactly what GatedDeltaNet::load did until it started going through linear_no_bias.
  • Still open: the GDN input projections are dequantized to bf16 so the grouped qkvz layout can be built by interleaving rows (~4 GB on this model). Projecting in_proj_qkv and in_proj_z separately and assembling the outputs would keep them quantized.

3b. Measured on spark and evo (2026-09-15)

Same box, same checkpoint, same harness (a streaming /v1/chat/completions client counting content and reasoning_content deltas; every prompt carries a fresh nonce so a prefix cache cannot report a prefill rate the engine never achieved).

spark (GB10, RadixArk/Qwen3.8-27B-NVFP4-BF16-LMHead), hanzo-engine with the native NVFP4 kernel:

prompt prefill, scalar tile prefill, WMMA decode
181 215.5 296.6 6.99
1,293 236.9 357.9 6.91
12,453 214.5 309.9 6.08

Tensor cores buy 38-51% of prefill and nothing on decode, which is the expected shape: a 16-wide MMA wastes 15 of its slots at batch 1. A wider K tile (two uint4 per row, halving the A-tile reloads and the barriers) measured 15-22% SLOWER -- 16 KB more shared memory costs more occupancy than the reuse buys -- and is recorded as a dead end in the kernel. Decode reads ~31 GB of weights per token against the checkpoint's computed 19.44 GB (below) and moves them at ~217 GB/s, which is 79% of the 273 GB/s the box has -- so decode is not slow, it is reading too much, and the ~4 GB is the GDN input projections this tree dequantizes. Prefill is the real gap and it is arithmetic: the MLP GEMMs run at roughly a tenth of the tensor core's rate.

spark after porting upstream's block-scaled FP4 GEMM (perf/nvfp4-prefill, d4c711170a). The CUTLASS kernel builds for sm_121a on CUDA 13.0; upstream's 13.3 pin and cuTile activation quantizer are not needed by the kernel, so the port quantizes activations with its own CUDA kernel against the checkpoint's input_scale. Decode (M <= 4) stays on the vecmat kernel.

engine prompt tokens prefill T/s decode T/s
hanzo-engine, FP4 GEMM 179 / 1,293 / 12,453 648 / 950 / 646 7.00 / 6.86 / 6.56
hanzo-engine, WMMA 175 / 1,293 / 12,453 297 / 358 / 310 6.99 / 6.91 / 6.08

Where the remaining prefill time goes (nsys, pp1291, GPU kernel time): bf16 GEMM over the FP8 attention layers that pertensor_fp8 dequantizes at load 27.3%, copies 13.2%, FP4 GEMM plus activation quantize 13.0%, elementwise 11.8%, casts 10.6%, GDN recurrence 9.2%. The same dequantized FP8 layers set decode: 26.65 GB read per token against the computed 19.44, at 187 GB/s. Native FP8 linear is the next lever for both.

The harness is checked in (scripts/bench_http.py, tests in scripts/bench_http_test.py). Prompt = uuid nonce + WORDS x size + "Summarize in one sentence."; sizes 4/40/400/4000 are ~220 / 1,335 / 12,497 / 124,094 Qwen3.8 tokens, and the prompt length reported is the server's usage.prompt_tokens. A token chunk is one whose delta carries content, reasoning, reasoning_content or tool_calls. TTFT and the decode window use token chunks only; decode = (usage.completion_tokens - 1) / (last token chunk - first token chunk), so MTP/DFlash steps that carry several tokens per chunk are counted as tokens, and a stream without usage fails (usage may arrive on the finish chunk or in a later empty-choices chunk). --concurrency N reports each stream plus the aggregate; --json FILE appends JSON lines.

python3 scripts/bench_http.py http://127.0.0.1:30000/v1/chat/completions --sizes 4,40,400,4000 --gen 128
python3 scripts/bench_http.py http://127.0.0.1:30000/v1/chat/completions --sizes 40 --concurrency 8

GB/token is two labeled numbers (scripts/bytes_per_token.py):

  • computed weight+KV+state bytes per token from the safetensors headers and config.json, skipping visual.* and mtp.* and counting one embed_tokens row; --resident REGEX=DTYPE names modules an engine holds at another dtype. weights --checkpoint <27B snapshot> --context 1400 --kv-dtype bf16 prints 19.44 GB of weights (19,435,407,488 B), 0.09 GB of bf16 KV (64 KiB/token over 16 full-attention layers) and 0.31 GB of GDN state: 19.84 GB per token.
  • measured SM-to-L2 bytes per token: sudo nsys profile --gpu-metrics-devices=0 --gpu-metrics-set=gb20y-top, then nsys export --type sqlite and bytes_per_token.py nsys --sqlite F --tokens N (sums 'VidL2 Total from L1TEX [Bytes]'). GB10's metric sets (gb20y, gb20y-top) have no DRAM counter (RmProfilingAdminOnly=1, so metrics need root), so this is L2 traffic including hits: a cross-check that exposes redundant reads, never a substitute for the computed figure.

evo (Strix Halo gfx1151, Qwen3.8-27B Q6_K), hanzo-engine ROCm:

hanzo-engine
prefill 512 51.2 T/s
prefill 4,096 44.3 T/s
decode 128 7.3 T/s

hanzo-engine bench --pa-context-len 16384; the serving harness agrees (35.8 T/s prefill / 5.15 decode). Prefill is the one to chase: gfx1151 has no tiled prefill path for Q6_K in our Vulkan backend and ROCm is running the eager attention.

Two things had to be fixed before any of this could be measured at all, and both are worth knowing:

  • The pager's demand floor ignored an explicit budget. On unified memory the KV cache is floored at one full model context, so --pa-memory-mb/--pa-context-len were raised back to 65 GB on a 124 GB box; a 27B bench then thrashed evo badly enough that SSH stopped answering and it took a privileged k3s pod on the node to kill the process. The floor now applies only to the automatic path.
  • A FlashInfer refusal used to kill the process. flashinfer_decode let flashinfer::Error cross back into Rust: "fatal runtime error: Rust cannot catch foreign exceptions". Qwen3.8-27B-NVFP4 aborted on its first token unless FLASHINFER_DECODE=0. Both decode entry points now return a status like the prefill one, and a refusal falls through to the standard paged kernel.
  • ROCm has no scatter/scatter_add in hanzo-ml, and interleaved MRoPE used Tensor::scatter, so every Qwen3.5 GGUF step failed with "scatter_set not yet implemented for ROCm". It is a gather now.
  • FlashInfer was chosen for a GQA group it cannot decode. The gate tested group <= 8, which is the prefill packing factor; the decode kernel instantiates 1, 2, 3, 4 and 8. Qwen3.5-27B has 6. Falling back at the call site cannot rescue it either -- the standard kernel reads block tables that are only built when it is the chosen backend -- so the gate is where it has to be right.
  • Long context on spark is blocked by the prefill attention, not by the quantization. With FlashInfer refused for this GQA group, prefill runs the eager path, which materializes the score matrix: a 4096-token chunk against a 124K prefix is 24 heads x 4096 x 124K x 4 B ~ 48 GB. The box wired ~110 of its 121 GB and stopped answering SSH. So 12.5K works and 124K does not, and the fix is a prefill that never builds the matrix (flash or FlashInfer), which is the "one prefill path" item in the plan.
  • A 124K prompt died in cumsum. hanzo-ml spells it as a triangular matmul, so the text-only MRoPE branch asked for a seq x seq matrix -- 15 billion elements -- and the launch came back CUDA_ERROR_INVALID_VALUE. The scan is linear and the mask is one byte per token, so it runs on the host now (positions_from_mask).
  • The GDN chunked prefill returned f32. Its fast path shadows q with its own f32 cast and then reads q.dtype() for the result, so any prompt of 64 tokens or more (CHUNK_THRESHOLD) came back f32 and met a bf16 weight in the next projection: "unexpected dtype, expected: F32, got: BF16" out of cublasLt. Under 64 tokens the same model answered fine, so it read as a streaming or sampling bug until the prompt length was swept.

4. Cloudflare Tunnel & Public Reselling Ingress (api.hanzo.ai)

  • Redundant cloudflared daemons run natively on host nodes and as Kubernetes pods (hanzo/cloudflared-*).
  • Ingress hanzo-domains-api-hanzo-ai terminates public traffic and routes to cloud:8000 (which routes through hanzo-router:1235 to Spark and Evo).
  • Exposes full OpenAI /v1/chat/completions and Anthropic /v1/messages (with structured thinking blocks).
  • Metered via Hanzo IAM: authenticated calls via Bearer <token> debit user/org quota and emit asynchronous usage records to the platform accounting database.

5. Workstation Setup & Claude Code Isolation (ra and dbc)

  • Standard claude remains completely clean and untouched: no global ANTHROPIC_BASE_URL or ANTHROPIC_API_KEY exports in shell profiles.
  • Dedicated launchers installed in ~/.local/bin/:
    • claude-hanzo picks the model from the name it is invoked by and execs claude; 1-hour stream timeouts; CLAUDE_CODE_AUTO_COMPACT_WINDOW=1000000. The window itself comes from the [1m] model suffix (see 2b), not from CLAUDE_CODE_MAX_CONTEXT_TOKENS, which is ignored for a model carrying a behavesAs row.
    • Effort is left to ~/.claude-hanzo/settings.json ("effortLevel": "medium"); exporting CLAUDE_CODE_EFFORT_LEVEL would pin it and make in-session /effort a no-op.
    • Defaults to https://api.hanzo.ai with the IAM token from hanzo auth token. --local uses the LAN router, which on macOS needs Local Network permission for ClaudeCode.app -- without it every LAN request fails as FailedToOpenSocket while curl from the same shell succeeds.
    • Dedicated configuration isolated to ~/.claude-hanzo/settings.json, preserving ~/.claude/settings.json for Anthropic/claude.ai.
  • Symlinks: claude-zen, claude-spark, claude-evo, claude-flash -> claude-hanzo. ~/.zshrc.d/hanzo-mesh.zsh holds no aliases or functions of those names: an alias adding --model zen5.8-coder drops the [1m] suffix and silently caps the window at 200K, and a same-named shell function shadows the executable and never launches claude.

6. Environment & Network Control (mainnet / testnet / devnet / local)

  • Switch active environment via hanzo network use <network>:
    • hanzo network use mainnet: Production L1 (chain 36963, RPC https://rpc.hanzo.network, API https://api.hanzo.ai).
    • hanzo network use testnet: Testnet (chain 36962, RPC https://rpc.testnet.hanzo.network).
    • hanzo network use devnet: Devnet (chain 36964, RPC https://rpc.devnet.hanzo.network).
    • hanzo network use local: Local development network (chain 1337).
  • Manage agent sandboxes and workspaces with hanzo sbx (hanzo sbx list, hanzo sbx run <agent>, hanzo sbx models).

7. Native Anthropic Server-Side Tools Support (web_search_*, bash_*, text_editor_*, computer_*)

  • Problem: When Claude Code issued web searches or tool calls, the request sent server-side tool definitions (e.g. type: "web_search_20250101", tool_choice: {"type": "tool", "name": "web_search"}). The replica's Anthropic adapter dropped these tools as unsupported, causing ValueError: tool_choice='tool' requires at least one custom tool; all supplied tools were server-side Anthropic built-ins which the OpenAI-compatible backend cannot invoke (HTTP 400).
  • Fix: spark's Anthropic adapter was patched to automatically synthesize function schemas for all Anthropic server-side built-in families:
    • web_search_* -> function web_search with {query: string} parameter schema.
    • bash_* -> function bash with {command: string, restart: boolean} schema.
    • text_editor_* -> function str_replace_editor / text_editor with full command parameters (view, create, str_replace, insert, undo_edit).
    • computer_* -> function computer with GUI action schemas (key, type, mouse_move, click, etc.).
  • Result: The underlying Qwen 3.8 / Zen model generates function calls normally; the adapter translates them back to Anthropic tool_use events (name: "web_search"), which Claude Code receives, executes, and continues with zero 400 errors.

8. Model Selection Consistency & Sub-Agent Routing on Evo

  • Consistent Model Naming & 1M Context UI:
    • zen5.8-coder: Zen 5.8 Coder (1M Context · DGX Spark) (claude-sonnet-4-6[1m], medium effort)
    • zen5.8-evo: Zen 5.8 Evo (1M Context · AMD Strix Halo) (claude-sonnet-4-6[1m], medium effort)
    • zen5-flash: Zen 5 Flash (Fast · Strix Halo Evo) (claude-haiku-4-5, subagents & auxiliary tasks)
    • zen5.8: Zen 5.8 (1M Context · DGX Spark) (claude-opus-5[1m])
  • Launcher Binary Auto-Selection:
    • /Users/z/.local/bin/claude-hanzo inspects basename $0:
      • claude-zen: automatically adds --model zen5.8-coder[1m].
      • claude-evo: automatically adds --model zen5.8-evo[1m].
      • claude-flash: automatically adds --model zen5-flash (no suffix; it is the short-context subagent model).
      • claude-spark: automatically adds --model zen5.8-spark[1m].
      • The picker rows in ~/.claude-hanzo/settings.json carry the same [1m] suffixes, so switching model in-session keeps the 1M window.
    • Automatically injects --dangerously-skip-permissions by default (overridable with --safe or --ask).
    • Configures CLAUDE_CONFIG_DIR=~/.claude-hanzo so standard claude (Claude Max / OAuth) remains 100% clean and untouched.
  • Sub-Agent Steering:
    • Sets CLAUDE_CODE_SUBAGENT_MODEL="zen5-flash", ANTHROPIC_DEFAULT_HAIKU_MODEL="zen5-flash", and ANTHROPIC_SMALL_MODEL="zen5-flash".
    • In router-pool.yaml, zen5-flash has weight 1000 on Strix Halo Evo APU, offloading fast subagents and summaries from Spark's 1M context main engine.

9. Hanzo Cloud MCP & Tooling Architecture

  • Automatic MCP Registration:
    • Registered hanzo-mcp across both ra and dbc workstations (/Users/z/.local/bin/hanzo-mcp and /Users/a/.local/bin/hanzo-mcp) in ~/.claude-hanzo/.claude.json.
    • Both machines verified healthy via CLAUDE_CONFIG_DIR=~/.claude-hanzo claude mcp list reporting hanzo: ... - ✔ Connected.
  • Local MCP Capabilities (31 Tools):
    • Exposes 31 native tools: browser, cdp, playwright (headless browser execution), fetch, curl, wget, git, code, fs, lsp, todo, vision, exec, ps, npx, uvx, open, jq, version, tool, and the unified dispatcher hanzo.
    • These tools run on the local host with zero remote authentication dependencies.
  • Unified Cloud MCP Projection (hanzo):
    • Rather than exposing all 120 cloud operations individually into the model's tool schema—which bloats the context and degrades tool routing accuracy—hanzo-mcp exposes a unified hanzo tool.
    • At runtime, hanzo dynamically introspects https://api.hanzo.ai/v1/openapi.json to route to all 120 services and operations:
      • Web Search: hanzo(service="websearch", action="search", params='{"query": "..."}')
      • Web Crawling: hanzo(service="crawl", action="...", params='...')
      • Cloud Infrastructure: hanzo(service="paas", ...), hanzo(service="iam", ...), hanzo(service="billing", ...), etc.
      • Catalog exploration: hanzo(service="services").
  • Direct Cloud MCP Endpoint (api.hanzo.ai/v1/mcp) Protocol Compatibility:
    • https://api.hanzo.ai/v1/mcp (in cloud/client/mcp.go) hardcodes const protocolVersion = "2026-07-28" in its initialize response.
    • Claude Code v2.1 client initiates MCP handshakes requesting 2025-11-25 and validates against its supported version array: ["2025-11-25", "2025-06-18", "2025-03-26", "2024-11-05", "2024-10-07"].
    • Because api.hanzo.ai unconditionally returns 2026-07-28 without down-negotiating to the client's requested protocol version, Claude Code throws: Server's protocol version is not supported: 2026-07-28.
    • hanzo-mcp operates as the local MCP bridge negotiating within Claude Code's supported protocol versions (2024-11-05 / 2025-11-25) while dispatching requests to api.hanzo.ai.
  • Authentication:
    • Authenticated cloud services (websearch, crawl, search) enforce tenant isolation and validated IAM principals (401: web search requires a validated principal).
    • Running hanzo auth login establishes/refreshes the OIDC session in ~/.hanzo/credentials.json, which hanzo-mcp automatically adopts.

MoE router and shared-expert overlap, held bitwise to the reference (W8c: E2, E3)

  • One router. src/cuda/exact/route.cu is a verbatim port of a reference top-k gating kernel (its header records the source and licence), built in the exact lane (src/cuda/exact/*.cu, no --use_fast_math). It serves every MoE model through ops::moe_router_topk for 15 fused expert counts (1-512 powers of two, 192, 320, 384, 448, 576); other counts keep the portable path. Bitwise vs the recorded reference outputs on 495 cases (15 counts x f32/bf16/f16 x 11 configs); the same source with fast math differs in 173,639 words.
  • The gate is cuBLAS's choice, and the workspace changes it. cuda::lane::Lane::linear is torch's F.linear (GemmEx T/N, COMPUTE_32F, DEFAULT_TENSOR_OP) on a handle carrying torch's 32 MiB workspace. Without it (hanzo-ml's handle, or 4 MiB) the N=512 gate differs at M 98-101, 126-160, 184-256, 289-320. Never route the gate through hanzo-ml matmul or UnquantLinear.
  • cuBLAS class map on GB10, cuBLAS 13.1.1 (rows bitwise equal across M within a class): N=512: {1, 17-183, 257-320} | {2-16} | {184-256, 321-4096}; N=1: one class. {2-16} is the reduced-precision split-K kernel (57-59% of logits correctly rounded vs 100% at M=1). The capture lists [1,2,4,8,16] and [1,2,4,5,8,10,16,20,...,80] never pad across a class; a full 512 grid would, at T=177..183 -> 184 (C3/W3 owns batch padding).
  • Pin. lane::CUBLAS = 13.1.1 (the reference was recorded on nvidia-cublas 13.1.1.3, byte-identical to /usr/local/cuda-13.0's). Hazard: /usr/local/cuda-13.4 with cuBLAS 13.7.0.27 is installed beside it; an alternatives flip or LD_LIBRARY_PATH change moves gate logits silently except for the startup warn, lane::tests::cublas_version and pad_classes. Dockerfile.cuda pins libcublas-13-0=13.1.1.3-1.
  • Branch (cuda::branch): side Lane + two reused events, fork before the router gate GEMM, join before the combine, capture-safe; refuses a context with cudarc event tracking on.
  • Goldens: hanzo-engine/tests/fixtures/zen6_moe/ holds the recorded reference router, gate and shared-expert outputs.