diff --git a/CHANGELOG.md b/CHANGELOG.md index f2a24e9..8a7bd6d 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -19,12 +19,16 @@ ### Added -- `--weights nvfp4|full` on `scripts/build-engine.sh` and `recipe/pilot.py`: the NVFP4 pruned weights serve with - the same two commands as the full model. The build checks out the engine pinned in `engine/nvfp4/UPSTREAM`; the - launcher downloads the pinned checkpoint on first use, checks it against `recipe/checksums/nvfp4.sha256` on every - start, refusing missing or changed files before it takes the GPU lock, and serves it with - `--cfg-gamma 1.5 --cfg-batched` unless a guidance flag of your own follows `--`. Without `--weights`, both behave - as before. +- Native Apple-silicon support for the pinned NVFP4-pruned checkpoint: packed experts, compact + original-ID routing and reference-conditioned semantic CFG, with independent backbone states. + The Mac engine uses `engine/nvfp4/macos/UPSTREAM` and its checksum-verified patches; the existing + full-model and CUDA engine pins are unchanged. [Mac qualification](results/macos-nvfp4.json) + records actual speech, streaming and guidance checks, including the initially failed stress fixture. +- `--weights nvfp4|full` on `scripts/build-engine.sh`, `recipe/pilot.py` and `recipe/pilot_macos.py`. + Selection verifies the pinned checkpoint on every start, including explicit checkpoint overrides, + and downloads only an absent default directory. Linux verifies before its GPU lock; Mac verifies + under its cooperative lock, before model spawn. NVFP4 defaults to `--cfg-gamma 1.5 --cfg-batched` + unless an explicit guidance flag follows `--`; legacy explicit-path mode retains operator-owned provenance. - Every speed or quality claim in a pull request carries before/after evidence from the recipe's bench and qualification scripts, measured against the recipe baseline (`CONTRIBUTING.md`, pull request template). - A speed claim also carries Grill v0.11.0 speech evidence: A, B and A2 captures of the recipe server and a diff --git a/NOTICE b/NOTICE index aaf8955..ab9ca99 100644 --- a/NOTICE +++ b/NOTICE @@ -4,19 +4,22 @@ Copyright 2026 plotarmordev License ------- All code in this repository is licensed under the Apache License, Version 2.0 (LICENSE). That covers the recipe, -scripts, benchmark and qualification tools, results, documentation and the engine patches in engine/patches/. +scripts, benchmark and qualification tools, results, documentation and the engine patches in engine/patches/ +and engine/nvfp4/macos/patches/. TensorFold ---------- The engine patches apply to TensorFold (https://github.com/ashhart/TensorFold), which is licensed under the Apache License, Version 2.0. scripts/build-engine.sh clones TensorFold at the commit pinned in engine/UPSTREAM -and applies the patches with `git am`. With --weights nvfp4 it clones instead the TensorFold fork -https://github.com/plotarmordev/TensorFold at the commit pinned in engine/nvfp4/UPSTREAM, also Apache-2.0, and -applies no patches. +and applies the patches with `git am`. On Linux, --weights nvfp4 instead clones the TensorFold fork +https://github.com/plotarmordev/TensorFold at the commit pinned in engine/nvfp4/UPSTREAM, also Apache-2.0, +and applies no patches. On Darwin, --weights nvfp4 uses upstream TensorFold at engine/nvfp4/macos/UPSTREAM +and applies the ordered patches in engine/nvfp4/macos/patches/. The patches modify TensorFold files. They add opt-in serving features and Drama 3 fixes: new files under -zig/src/server/, zig/src/core/, zig/src/native/, docs/ and tools/zig/, and changes to existing TensorFold files in -zig/build/, zig/src/core/, zig/src/families/drama3/, zig/src/native/ and zig/src/server/, plus README.md and +zig/src/server/, zig/src/core/, zig/src/native/, zig/kernels/metal/, zig/tests/, docs/ and tools/zig/, +and changes to existing TensorFold files in zig/build/, zig/src/core/, zig/src/families/drama3/, +zig/src/native/ and zig/src/server/, plus README.md and CHANGELOG.md. Each patch's commit message and diff state what it changes. The patches keep TensorFold's existing file headers and notices unchanged. diff --git a/README.md b/README.md index f500110..589d3f1 100644 --- a/README.md +++ b/README.md @@ -22,12 +22,12 @@ > [Fish Audio Drama 3](https://huggingface.co/fishaudio/drama-3) under the Fish Audio Community License v1.0 > ([License](#license)). `--weights full` serves the full model, the alternative. -**Fish Audio's Drama 3 text-to-speech on an NVIDIA GB10, on TensorFold's native Zig/CUDA engine.** +**Fish Audio's Drama 3 text-to-speech on NVIDIA GB10 (native Zig/CUDA) or Apple silicon (native Zig/Metal).** [TensorFold](https://github.com/ashhart/TensorFold) runs Drama 3 natively since v1.1.0. This repository pins -upstream TensorFold and adds a small series of opt-in patches for the full model; for the NVFP4 weights it pins -the TensorFold engine they were measured on. It builds the engine and starts a local speech server with an -OpenAI-style `/v1/audio/speech` route, whole or streamed. +upstream TensorFold and adds a small series of opt-in patches for the full model. NVFP4 retains its +measured CUDA engine pin and uses a separate upstream pin plus Metal patches on Mac. It builds the engine +and starts a local speech server with an OpenAI-style `/v1/audio/speech` route, whole or streamed. **Built with Fish Audio.** The model weights are not in this repository; the launcher downloads them from [Hugging Face](https://huggingface.co/plotarmordev/drama-3-NVFP4-Pruned) under Fish Audio's own @@ -40,9 +40,9 @@ OpenAI-style `/v1/audio/speech` route, whole or streamed. ## Quickstart -Apple silicon: use the [full-model macOS/Metal instructions](recipe/README-macos.md). The NVFP4 -comparison is on GB10; it does not qualify that engine on Mac. The Linux requirements and -cgroup cap below do not apply to the Mac launcher. +Apple silicon: use the [macOS/Metal instructions](recipe/README-macos.md) for `--weights nvfp4` +or `--weights full`. The GB10 measurements above do not qualify Mac speed or quality. +The Linux requirements and cgroup cap below do not apply to the Mac launcher. You need Linux with user systemd, CUDA 13, [Zig 0.17.0](https://ziglang.org/download/), the `hf` command (`python3 -m pip install -U huggingface_hub`) and 88 GiB of free memory at load (the server is capped at 54 GiB). @@ -64,10 +64,11 @@ For the full model, write `--weights full` in both commands. The first start downloads the weights at the revision pinned in [`recipe/sources.json`](recipe/sources.json), into `drama-3-NVFP4-Pruned/` or `drama-3/`. Every start checks them against [`recipe/checksums`](recipe/checksums) and refuses missing or changed files; for the 21 GiB NVFP4 weights that takes about 11 seconds on a GB10. -With `--weights nvfp4`, `scripts/build-engine.sh` builds the engine pinned in -[`engine/nvfp4/UPSTREAM`](engine/nvfp4/UPSTREAM) and the launcher serves with guidance γ 1.5 and batched guidance, -as measured; a `--cfg-gamma` of your own after `--` replaces both. With `--weights full` it checks out the commit -pinned in [`engine/UPSTREAM`](engine/UPSTREAM) and applies [`engine/patches`](engine/patches) in order. +On Linux, `--weights nvfp4` builds [`engine/nvfp4/UPSTREAM`](engine/nvfp4/UPSTREAM) and serves with +guidance γ 1.5 and batched guidance, as measured; a `--cfg-gamma` of your own after `--` replaces both. +On Mac it builds [`engine/nvfp4/macos/UPSTREAM`](engine/nvfp4/macos/UPSTREAM) plus its patch series. +Metal's `--cfg-batched` pairs independent states in one command buffer, not a vectorized GEMM. +`--weights full` retains [`engine/UPSTREAM`](engine/UPSTREAM) and [`engine/patches`](engine/patches). When it prints `DRAMA_PILOT_READY`, synthesize from another terminal. The second command streams and plays as it arrives: @@ -122,9 +123,10 @@ the license's Section III(c), which needs a separate license from Fish Audio. `POST /v1/audio/speech` takes `model` (`fishaudio/drama-3`), `input` (at most 4096 characters), `voice` (`default`, or a name loaded with `--voices`), `response_format` (`wav` or `pcm`) and `stream` (`true` needs `pcm`). Request-supplied references and other synthesis modifiers are refused. `GET /v1/voices` lists the voices. -With `--weights nvfp4`, `--voices` takes the `drama3-voice` files its encoder writes, a request may also send its -own `references` as codec codes (audio is refused), and every reply carries the server's γ in an `X-Cfg-Gamma` -header, though only a request with a voice is drawn with guidance. +The Linux CUDA NVFP4 engine accepts its encoder's `drama3-voice` catalogs and request `references` +as codec codes (not audio). The Mac NVFP4 engine instead uses attested `drama3-http-voice/1` +catalogs and refuses request-supplied references; see the [Mac voice commands](recipe/README-macos.md#attested-reference-voices). +Guidance-enabled replies carry the configured γ in `X-Cfg-Gamma`; builtin voices remain unguided. - **Whole reply:** `Content-Length` framing and an `X-Finish-Reason` header: `stop`, `length` or, with the silence guard on, `silence`. diff --git a/engine/nvfp4/macos/UPSTREAM b/engine/nvfp4/macos/UPSTREAM new file mode 100644 index 0000000..71ff469 --- /dev/null +++ b/engine/nvfp4/macos/UPSTREAM @@ -0,0 +1,4 @@ +# Native Metal base for the NVFP4 checkpoint; patches/ adds its packed experts, compaction and guidance. +# The CUDA NVFP4 pin at ../UPSTREAM and the full-model pin at ../../UPSTREAM remain unchanged. +TENSORFOLD_REPOSITORY=https://github.com/ashhart/TensorFold.git +TENSORFOLD_COMMIT=3d219d2156d4c25dcfb74370c8703cd89e58aa49 diff --git a/engine/nvfp4/macos/patches/0001-drama3-metal-nvfp4-cfg.patch b/engine/nvfp4/macos/patches/0001-drama3-metal-nvfp4-cfg.patch new file mode 100644 index 0000000..0f49368 --- /dev/null +++ b/engine/nvfp4/macos/patches/0001-drama3-metal-nvfp4-cfg.patch @@ -0,0 +1,3395 @@ +From 5d6f8a09c6091ffba2a7ef65f43eb503a4ff4ca9 Mon Sep 17 00:00:00 2001 +From: jerryfane <38183696+jerryfane@users.noreply.github.com> +Date: Sun, 25 Oct 2026 04:57:47 +0400 +Subject: [PATCH] drama: support NVFP4 and reference CFG on Metal + +Keep expert storage packed and retain checkpoint activation precision while +mapping original router IDs to compact expert slots. Independent backbone +states allow reference-conditioned semantic guidance without changing codec +heads or builtin-voice behavior. + +On Apple M5 Ultra, the production code completed 28 HTTP checks: 15 complete +PCM replies, one early stream close and 12 expected validation refusals. +Paired and separate guidance replies matched exactly. The initial registered +operator run passed 13/14 variants and all runtime limits; the following +test-only patch corrects the remaining fixture's precision discrimination. + +--- + CHANGELOG.md | 3 + + README.md | 5 + + zig/build/cuda.zig | 2 + + zig/build/drama3_host.zig | 5 + + zig/build/drama3_metal.zig | 13 +- + zig/build/kernel_contract.zig | 8 + + zig/kernels/metal/drama3/guidance.metal | 14 + + zig/kernels/metal/drama3/nvfp4.metal | 135 +++++++++ + zig/kernels/metal/sources.zig | 2 + + zig/src/core/contract.zig | 1 + + zig/src/core/contract/registry.zig | 5 + + zig/src/core/engine_api.zig | 4 + + zig/src/core/speech_api.zig | 3 + + zig/src/families/drama3/drama3_metal.zig | 4 + + zig/src/families/drama3/metal/engine.zig | 49 ++- + zig/src/families/drama3/metal/forward.zig | 36 ++- + zig/src/families/drama3/metal/gpu.zig | 5 + + zig/src/families/drama3/metal/guidance.zig | 152 ++++++++++ + zig/src/families/drama3/metal/kernels.zig | 20 +- + zig/src/families/drama3/metal/load_limit.zig | 16 +- + zig/src/families/drama3/metal/quant_check.zig | 232 ++++++++++++++ + .../families/drama3/metal/quant_format.zig | 73 +++++ + .../drama3/metal/quant_format_test.zig | 90 ++++++ + .../families/drama3/metal/quant_reference.zig | 63 ++++ + zig/src/families/drama3/metal/weights.zig | 79 +++-- + zig/src/families/drama3/shards.zig | 2 + + zig/src/native/drama_host.zig | 41 ++- + zig/src/native/drama_open.zig | 6 + + zig/src/native/metal.zig | 4 + + zig/src/server/cli.zig | 10 +- + zig/src/server/drama_cfg.zig | 121 ++++++++ + zig/src/server/drama_prompt.zig | 10 + + zig/src/server/root.zig | 2 + + zig/src/server/serve.zig | 1 + + zig/src/server/server.zig | 1 + + zig/src/server/speech.zig | 25 +- + zig/src/server/speech_cfg_test.zig | 59 ++++ + zig/src/server/speech_metadata.zig | 5 + + zig/src/server/speech_redraw_test.zig | 1 - + zig/src/server/speech_stream_test.zig | 5 + + zig/tests/contract/drama3_nvfp4_ref.zig | 223 ++++++++++++++ + zig/tests/contract/metal_drama3.zig | 2 +- + zig/tests/contract/metal_drama3_nvfp4.zig | 239 +++++++++++++++ + zig/tests/contract/metal_drama_cfg.zig | 154 ++++++++++ + zig/tests/contract/metal_drama_combine.zig | 285 ++++++++++++++++++ + zig/tests/contract/metal_drama_router.zig | 224 ++++++++++++++ + zig/tests/contract/metal_main.zig | 40 +++ + zig/tests/drama_nvfp4_metal.zig | 10 + + 48 files changed, 2418 insertions(+), 71 deletions(-) + create mode 100644 zig/kernels/metal/drama3/guidance.metal + create mode 100644 zig/kernels/metal/drama3/nvfp4.metal + create mode 100644 zig/src/families/drama3/metal/guidance.zig + create mode 100644 zig/src/families/drama3/metal/quant_check.zig + create mode 100644 zig/src/families/drama3/metal/quant_format.zig + create mode 100644 zig/src/families/drama3/metal/quant_format_test.zig + create mode 100644 zig/src/families/drama3/metal/quant_reference.zig + create mode 100644 zig/src/server/drama_cfg.zig + create mode 100644 zig/src/server/speech_cfg_test.zig + create mode 100644 zig/tests/contract/drama3_nvfp4_ref.zig + create mode 100644 zig/tests/contract/metal_drama3_nvfp4.zig + create mode 100644 zig/tests/contract/metal_drama_cfg.zig + create mode 100644 zig/tests/contract/metal_drama_combine.zig + create mode 100644 zig/tests/contract/metal_drama_router.zig + create mode 100644 zig/tests/drama_nvfp4_metal.zig + +diff --git a/CHANGELOG.md b/CHANGELOG.md +index 41d25ebd..50a2c7ad 100644 +--- a/CHANGELOG.md ++++ b/CHANGELOG.md +@@ -4,6 +4,9 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel + + ## Unreleased + ++- Drama 3 on Metal reads the pruned NVFP4 checkpoint's packed E2M1 experts, block-16 E4M3 scales and F32 global scales without requantizing or expanding the experts. Routing masks dropped original IDs before softmax and maps selected IDs to compact slots. The full FP8-expert path retains its existing arithmetic. ++- Metal Drama speech supports opt-in semantic classifier-free guidance with shared weights and independent conditional/unconditional backbone state. `--cfg-gamma`, `--cfg-allow-request-gamma`, `--cfg-force-guided` and `--cfg-batched` preserve builtin unguided voices, validate request gamma, and reset guidance on cancellation or reference changes. Paired mode submits both states in one command buffer, not a vectorized GEMM; admission includes the extra state. CUDA guidance is not enabled by this change. ++ + - Drama speech on Metal reuses a catalog voice's reference prefill across requests on the same lane, so a request with the same voice prefills only its own text. Audio is byte-identical to an uncached request. On an M3 Ultra with Drama 3 and a catalog voice, the first audio byte of a streamed reply comes 90% sooner (15.6 s to 1.5 s over four texts) and whole replies take about half as long; the default voice and the first request with a voice are unchanged (#1237). + + - Flash Next on Macs without tensor units decodes faster with the same tokens: the hyper-connection act pair runs for windows of two rows and more, the HC down projection runs without staging or barriers at four outputs a threadgroup, rounds are gated on the GPU with a GPU-only event between them, drafts are gated at tau 0.4, the decode and draft argmaxes take two launches, and wide windows use grouped experts and exact tree windows. On an M3 Ultra with Qwen3.8 Flash Next MLX 4-bit MTP, the served sweep goes from 120.2 to 147.6 tokens/s on one stream (+22.7%), 129.7 to 155.3 at two (+19.8%), 143.8 to 167.0 at four (+16.2%) and 154.4 to 176.1 at eight (+14.1%); prefill is unchanged and every token is identical. +diff --git a/README.md b/README.md +index 795ffe2b..cf69ddec 100644 +--- a/README.md ++++ b/README.md +@@ -179,6 +179,7 @@ Conversions are explicitly marked: their upstream weights need the linked conver + | Kolibri-1, block FP8 (research) | `tensorfold pull Aleph-Alpha/Kolibri-1@e52eb4627d11516b0c01de49210ab5a4e4061444` | One NVIDIA GPU; GB10 measured; 74 GB checkpoint. | [Kolibri](docs/recipes/kolibri-1.md). Text only; copy drafts, no EAGLE-3/neural drafter. | + | Liquid AI d1-3B, decisions | Not documented yet. | Native CUDA; RAM minimum not published. | [API section below](#serve-flags). `/v1/systemone` and `/v1/decisions`. | + | Fish Audio Drama 3, speech | Not documented yet. | Native Metal or CUDA; GB10 measured. | [Release audio instructions](CHANGELOG.md#110). `/v1/audio/speech`; checkpoint and `codec.pth` required. | ++| Drama 3 pruned NVFP4, speech | `hf download plotarmordev/drama-3-NVFP4-Pruned --revision 494db08481db952f834b2241bce63fb3dfa72554 --local-dir drama-3-NVFP4-Pruned` | Native Metal; roughly 21 GiB checkpoint plus runtime state. Not a minimum-RAM or cross-backend quality claim. | Packed experts, BF16 dense/tower weights and FP32 accumulation; no conversion. Use `serve drama-3-NVFP4-Pruned --cfg-gamma 1.5 --cfg-batched` with an attested catalog voice. | + + ### GLM-5.3: one machine or two? + +@@ -287,6 +288,10 @@ The binary's `capabilities --json` response lists its supported flags and platfo + | `--drafter-bits 0\|4` | Drafter weight precision where the model offers a choice: 4 bits (the default), or 0 to keep bf16. Qwen3.8-27B on CUDA takes 4 only. | + | `--record DIR` | Write each stream's committed tokens, positions and tapped rows under DIR, for drafter training. CUDA Nemotron only; a single stream runs through the lane rounds, and recording stops below 25 GiB free disk. | + | `--keep-warm SECONDS` | Keep the Metal GPU active while idle for this long after a request, default 900; zero disables it. | ++| `--cfg-gamma GAMMA` | Metal Drama speech: semantic guidance strength, finite 1–5; 0/default disables it. Only reference-conditioned requests are guided. Enabled replies report `X-Cfg-Gamma`; codec heads stay conditional. | ++| `--cfg-allow-request-gamma` | Metal Drama speech: allow a finite 1–5 `X-Cfg-Gamma` request header; invalid values return 400 before engine submission. Without this opt-in, the server's gamma is fixed. | ++| `--cfg-force-guided` | With an enabled Metal Drama guidance policy, exercise the unconditional state even at gamma 1. The semantic mix at 1 remains exact conditional-logit identity. | ++| `--cfg-batched` | With enabled Metal Drama guidance, submit both independent backbone states in one command buffer rather than separate submissions. Weights are shared; extra state is included in admission. Not vectorized batching. | + | `--silence-floor-dbfs DBFS` | Drama speech only: explicit frame-power floor for the silence guard, with no default floor. Calibrate below quiet intended speech using public or built-in speech. | + | `--silence-seconds SECONDS` | Drama speech only: stop after this continuous duration below the explicit floor, retaining audio and reporting `silence`. Unset by default: off; calibrate above intentional pauses. | + | `--silence-onset-s SECONDS` | Drama speech only: with the silence guard on, stop once this many seconds run below its floor while the synthesis has voiced less than `--silence-onset-voiced-s`, retaining audio and reporting `silence`. Off unless given; must be below `--silence-seconds`. | +diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig +index c696b525..5f9393ac 100644 +--- a/zig/build/cuda.zig ++++ b/zig/build/cuda.zig +@@ -419,6 +419,8 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st + .imports = &.{ .{ .name = "engine_api", .module = native_modules.api }, .{ .name = "tokenizer", .module = mods.tokenizer }, .{ .name = "template", .module = template }, .{ .name = "living", .module = mods.living } }, + }); + const server_json = @import("server.zig").addTo(b, server_tests, host, .debug); ++ b.step("test-drama-cfg-host", "Drama CFG CLI, HTTP and voice transitions without GPU execution").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = server_tests, .filters = &.{"CFG"} })).step); ++ b.step("test-speech-redraw-deadline-host", "Busy speech redraw expiry without counting an unsubmitted take").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = server_tests, .filters = &.{"a busy redraw stops at its latest start"} })).step); + @import("proposal_vocab_host.zig").steps(b, mods.core); + const eos_policy = b.step("test-eos-policy-host", "Shared EOS, real HfText and HTTP policy without a GPU"); + eos_policy.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = server_tests, .filters = &.{"EOS policy:"} })).step); +diff --git a/zig/build/drama3_host.zig b/zig/build/drama3_host.zig +index 48e6b938..91088594 100644 +--- a/zig/build/drama3_host.zig ++++ b/zig/build/drama3_host.zig +@@ -46,6 +46,11 @@ pub fn metal(b: *std.Build, core: *std.Build.Module, api: *std.Build.Module) voi + const step = b.step("check-drama-metal-host", "Compile Drama's Metal host entry points without compiling shaders"); + compile(b, step, family, "drama_metal_host"); + native(b, step, family, core, api, null, "drama_metal_native_host"); ++ const quant = b.addRunArtifact(b.addTest(.{ ++ .root_module = family, ++ .filters = &.{ "NVFP4", "expert metadata", "independent scalar reference", "CFG" }, ++ })); ++ b.step("test-drama-metal-quant-host", "Packed expert formats, pruning and CFG state without GPU execution").dependOn(&quant.step); + } + + /// Drama's prompt plan on the host: the single-row default and the opt-in batched plan. +diff --git a/zig/build/drama3_metal.zig b/zig/build/drama3_metal.zig +index 7932179e..a74e0270 100644 +--- a/zig/build/drama3_metal.zig ++++ b/zig/build/drama3_metal.zig +@@ -2,7 +2,7 @@ + const std = @import("std"); + + const dir = "zig/kernels/metal/drama3"; +-const sources = [_][]const u8{ "dense.metal", "gdn.metal", "attn.metal", "moe.metal", "sample.metal", "codec.metal" }; ++const sources = [_][]const u8{ "dense.metal", "gdn.metal", "attn.metal", "moe.metal", "nvfp4.metal", "guidance.metal", "sample.metal", "codec.metal" }; + + /// Drama in the Metal engines: its `core` is the engine's safetensors, the server's tokenizer and frame_rounds. + pub fn addNative(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, metal: *std.Build.Module, native: *std.Build.Module, metal_std: []const u8) void { +@@ -65,6 +65,17 @@ pub fn targets(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.bu + .link_libc = true, + .imports = &.{ .{ .name = "metal", .module = metal }, .{ .name = "core", .module = core }, .{ .name = "drama3_metallib", .module = bytes } }, + }); ++ const quant_check = b.addExecutable(.{ ++ .name = "drama3-metal-quant-check", ++ .root_module = b.createModule(.{ ++ .root_source_file = b.path("zig/tests/drama_nvfp4_metal.zig"), ++ .target = target, ++ .optimize = optimize, ++ .link_libc = true, ++ .imports = &.{ .{ .name = "drama3", .module = family }, .{ .name = "drama3_metallib", .module = bytes } }, ++ }), ++ }); ++ b.step("drama3-metal-quant-check", "Build the NVFP4 Metal operator checker; does not execute GPU work").dependOn(&b.addInstallArtifact(quant_check, .{}).step); + const tool = b.createModule(.{ + .root_source_file = b.path("zig/src/families/drama3/generate.zig"), + .target = target, +diff --git a/zig/build/kernel_contract.zig b/zig/build/kernel_contract.zig +index 0a60668b..fcac3fff 100644 +--- a/zig/build/kernel_contract.zig ++++ b/zig/build/kernel_contract.zig +@@ -36,6 +36,14 @@ pub fn hostTests(b: *std.Build, family: *std.Build.Step) void { + }) })); + b.step("test-delta-slots-contract-host", "FP64 DeltaNet and slot corruption controls without shaders").dependOn(&delta.step); + step.dependOn(&delta.step); ++ const nvfp4 = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ ++ .root_source_file = b.path("zig/tests/contract/drama3_nvfp4_ref.zig"), ++ .target = b.graph.host, ++ .optimize = .debug, ++ .imports = &.{.{ .name = "contract", .module = module(b, b.graph.host, .debug) }}, ++ }) })); ++ b.step("test-drama-nvfp4-reference-host", "NVFP4 scalar reference and reduced-accumulator discrimination").dependOn(&nvfp4.step); ++ step.dependOn(&nvfp4.step); + } + + /// macOS: tf-kernel-contract over the engine's Metal kernels; `kernel-contract` runs it on this Mac's GPU. +diff --git a/zig/kernels/metal/drama3/guidance.metal b/zig/kernels/metal/drama3/guidance.metal +new file mode 100644 +index 00000000..78117694 +--- /dev/null ++++ b/zig/kernels/metal/drama3/guidance.metal +@@ -0,0 +1,14 @@ ++#include ++using namespace metal; ++ ++struct D3CfgArgs { uint count; float gamma; }; ++ ++// Candidate logits include semantic tokens and end-of-speech; residual codec heads remain conditional. ++[[max_total_threads_per_threadgroup(256)]] kernel void d3_cfg_mix(device float* conditional [[buffer(0)]], ++ device const float* unconditional [[buffer(1)]], ++ constant D3CfgArgs& args [[buffer(2)]], ++ uint i [[thread_position_in_grid]]) { ++ if (i >= args.count || args.gamma == 1.0f) return; ++ float c = conditional[i]; ++ conditional[i] = c + (args.gamma - 1.0f) * (c - unconditional[i]); ++} +diff --git a/zig/kernels/metal/drama3/nvfp4.metal b/zig/kernels/metal/drama3/nvfp4.metal +new file mode 100644 +index 00000000..451869bf +--- /dev/null ++++ b/zig/kernels/metal/drama3/nvfp4.metal +@@ -0,0 +1,135 @@ ++// Packed row-major NVFP4: E2M1 nibbles, unsigned E4M3 block16 scales, FP32 projection globals. ++#include "d3_common.h" ++ ++struct Fp4Projection { device const uchar* w; device const uchar* s; device const float* global; }; ++struct Fp4Expert { Fp4Projection g, u, d; }; ++struct Fp4Args { uint K, N, prompt, cap, slots; }; ++ ++inline float fp4_value(uchar code) { ++ const float values[8] = {0.0f, 0.5f, 1.0f, 1.5f, 2.0f, 3.0f, 4.0f, 6.0f}; ++ return (code & 8) ? -values[code & 7] : values[code & 7]; ++} ++inline float fp4_scale(uchar code) { ++ const uint exp = (code >> 3) & 15, mant = code & 7; ++ return exp ? ldexp(1.0f + float(mant) / 8.0f, int(exp) - 7) : float(mant) / 512.0f; ++} ++inline float fp4_row(Fp4Projection w, uint n, uint K, device const bf16* x, uint lane) { ++ float acc = 0.0f; ++ for (uint k = lane * 16; k < K; k += 512) { ++ float block = 0.0f; ++ for (uint j = 0; j < 16; ++j) { ++ const uchar b = w.w[(ulong)n * (K / 2) + (k + j) / 2]; ++ const uchar code = (j & 1) ? (b >> 4) : (b & 15); ++ block = fma(fp4_value(code), b2f(x[k + j]), block); ++ } ++ acc = fma(block, fp4_scale(w.s[(ulong)n * (K / 16) + k / 16]), acc); ++ } ++ return simd_sum(acc) * w.global[0]; ++} ++ ++[[max_total_threads_per_threadgroup(128)]] kernel void d3_nvfp4_up( ++ device const bf16* x [[buffer(0)]], device const Fp4Expert* table [[buffer(1)]], ++ device const int* picks [[buffer(2)]], device const int* counts [[buffer(3)]], ++ device const int* members [[buffer(4)]], device bf16* out [[buffer(5)]], ++ constant Fp4Args& a [[buffer(6)]], uint2 tg [[threadgroup_position_in_grid]], ++ uint sg [[simdgroup_index_in_threadgroup]], uint lane [[thread_index_in_simdgroup]]) { ++ const uint n = tg.x * 4 + sg; ++ const int e = a.prompt ? int(tg.y) : picks[tg.y]; ++ const int cnt = a.prompt ? counts[e] : 1; ++ if (!cnt) return; ++ const Fp4Expert w = table[e]; ++ for (int j = 0; j < cnt; ++j) { ++ const uint pair = a.prompt ? uint(members[(ulong)e * a.cap + j]) : tg.y; ++ device const bf16* row = x + (ulong)(pair / a.slots) * a.K; ++ const float g = rb(fp4_row(w.g, n, a.K, row, lane)); ++ const float u = rb(fp4_row(w.u, n, a.K, row, lane)); ++ if (lane == 0) out[(ulong)pair * a.N + n] = f2b(g / (1.0f + precise::exp(-g)) * u); ++ } ++} ++ ++[[max_total_threads_per_threadgroup(128)]] kernel void d3_nvfp4_down( ++ device const bf16* x [[buffer(0)]], device const Fp4Expert* table [[buffer(1)]], ++ device const int* picks [[buffer(2)]], device const int* counts [[buffer(3)]], ++ device const int* members [[buffer(4)]], device float* out [[buffer(5)]], ++ constant Fp4Args& a [[buffer(6)]], uint2 tg [[threadgroup_position_in_grid]], ++ uint sg [[simdgroup_index_in_threadgroup]], uint lane [[thread_index_in_simdgroup]]) { ++ const uint n = tg.x * 4 + sg; ++ const int e = a.prompt ? int(tg.y) : picks[tg.y]; ++ const int cnt = a.prompt ? counts[e] : 1; ++ if (!cnt) return; ++ const Fp4Expert w = table[e]; ++ for (int j = 0; j < cnt; ++j) { ++ const uint pair = a.prompt ? uint(members[(ulong)e * a.cap + j]) : tg.y; ++ const float v = fp4_row(w.d, n, a.K, x + (ulong)pair * a.K, lane); ++ if (lane == 0) out[(ulong)pair * a.N + n] = v; ++ } ++} ++ ++struct CompactRouterArgs { uint ss, off; }; ++// Retained original IDs compete before softmax/top8; only the selected IDs are translated to compact slots. ++[[max_total_threads_per_threadgroup(256)]] kernel void d3_router_compact( ++ device const bf16* sp [[buffer(0)]], device int* picks [[buffer(1)]], ++ device float* weights [[buffer(2)]], device const int* slots [[buffer(3)]], ++ constant CompactRouterArgs& a [[buffer(4)]], uint t [[threadgroup_position_in_grid]], ++ uint e [[thread_index_in_threadgroup]], uint sg [[simdgroup_index_in_threadgroup]], ++ uint lane [[thread_index_in_simdgroup]]) { ++ threadgroup float scratch[32], wv[8], best_v[8]; ++ threadgroup int wi[8], best_i[8]; ++ const bool kept = slots[e] >= 0; ++ const float lg = kept ? b2f(sp[(ulong)t * a.ss + a.off + e]) : -INFINITY; ++ const float m = group_max(lg, scratch, sg, lane, 8); ++ const float p = precise::exp(lg - m); ++ const float sum = group_sum(p, scratch + 8, sg, lane, 8); ++ // Poisoned probabilities must propagate NaNs without indexing a dropped expert or selecting one twice. ++ const bool poisoned = !isfinite(sum) || sum <= 0.0f; ++ float mine = kept ? (poisoned ? 0.0f : p / sum) : -1.0f; ++ for (int k = 0; k < 8; ++k) { ++ float v = mine; ++ int idx = int(e); ++ for (ushort o = 16; o > 0; o >>= 1) { ++ const float ov = simd_shuffle_xor(v, o); ++ const int oi = simd_shuffle_xor(idx, o); ++ if (ov > v || (ov == v && oi < idx)) { v = ov; idx = oi; } ++ } ++ if (lane == 0) { wv[sg] = v; wi[sg] = idx; } ++ threadgroup_barrier(mem_flags::mem_threadgroup); ++ if (e == 0) { ++ float bv = wv[0]; ++ int bi = wi[0]; ++ for (int j = 1; j < 8; ++j) ++ if (wv[j] > bv || (wv[j] == bv && wi[j] < bi)) { bv = wv[j]; bi = wi[j]; } ++ best_v[k] = bv; ++ best_i[k] = bi; ++ } ++ threadgroup_barrier(mem_flags::mem_threadgroup); ++ if (int(e) == best_i[k]) mine = -1.0f; ++ threadgroup_barrier(mem_flags::mem_threadgroup); ++ } ++ if (e < 8) { ++ float total = 0.0f; ++ for (int j = 0; j < 8; ++j) total += best_v[j]; ++ picks[t * 8 + e] = slots[best_i[e]]; ++ weights[t * 8 + e] = poisoned ? NAN : best_v[e] / total; ++ } ++} ++ ++// CUDA drama_moe_combine_f32: round each weighted pair, sum, then fuse the FP32 shared gate and add residual. ++[[max_total_threads_per_threadgroup(256)]] kernel void d3_nvfp4_combine( ++ device bf16* residual [[buffer(0)]], device const float* down [[buffer(1)]], ++ device const float* weights [[buffer(2)]], device const bf16* shared [[buffer(3)]], ++ device const bf16* x [[buffer(4)]], device const bf16* wg [[buffer(5)]], ++ constant uint& D [[buffer(6)]], uint row [[threadgroup_position_in_grid]], ++ uint t [[thread_index_in_threadgroup]], uint sg [[simdgroup_index_in_threadgroup]], ++ uint lane [[thread_index_in_simdgroup]]) { ++ threadgroup float scratch[32]; ++ float part = 0.0f; ++ for (uint i = t; i < D; i += 256) part += b2f(x[(ulong)row * D + i]) * b2f(wg[i]); ++ const float gate = d3_sigmoid(group_sum(part, scratch, sg, lane, 8)); ++ for (uint i = t; i < D; i += 256) { ++ float sum = 0.0f; ++ for (uint slot = 0; slot < 8; ++slot) ++ sum += rb(weights[row * 8 + slot] * down[((ulong)row * 8 + slot) * D + i]); ++ const float moe = rb(gate * b2f(shared[(ulong)row * D + i]) + rb(sum)); ++ residual[(ulong)row * D + i] = f2b(b2f(residual[(ulong)row * D + i]) + moe); ++ } ++} +diff --git a/zig/kernels/metal/sources.zig b/zig/kernels/metal/sources.zig +index 89047df0..b71562cb 100644 +--- a/zig/kernels/metal/sources.zig ++++ b/zig/kernels/metal/sources.zig +@@ -178,6 +178,8 @@ pub const drama3 = struct { + pub const gdn = @embedFile("drama3/gdn.metal"); + pub const attn = @embedFile("drama3/attn.metal"); + pub const moe = @embedFile("drama3/moe.metal"); ++ pub const nvfp4 = @embedFile("drama3/nvfp4.metal"); ++ pub const guidance = @embedFile("drama3/guidance.metal"); + pub const sample = @embedFile("drama3/sample.metal"); + pub const codec = @embedFile("drama3/codec.metal"); + }; +diff --git a/zig/src/core/contract.zig b/zig/src/core/contract.zig +index 4d427aec..1d5aff87 100644 +--- a/zig/src/core/contract.zig ++++ b/zig/src/core/contract.zig +@@ -6,6 +6,7 @@ pub const gates = @import("contract/gates.zig"); + pub const rows = @import("contract/rows.zig"); + pub const affine = @import("contract/affine.zig"); + pub const dense = @import("contract/dense.zig"); ++pub const bf16_round = @import("bf16_round.zig"); + pub const limits = @import("contract/limits.zig"); + pub const registry = @import("contract/registry.zig"); + pub const cubin = @import("contract/cubin.zig"); +diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig +index d5cb7cbe..c7e8f24f 100644 +--- a/zig/src/core/contract/registry.zig ++++ b/zig/src/core/contract/registry.zig +@@ -508,6 +508,11 @@ pub const metal = [_]M{ + .{ .name = "drama3 moe d3_plan", .threads = 1024, .tg_memory = 1024, .reserved = 1024 }, + .{ .name = "drama3 moe d3_expert_up", .threads = 128, .tg_memory = 0, .reserved = 128 }, + .{ .name = "drama3 moe d3_expert_down", .threads = 128, .tg_memory = 0, .reserved = 128 }, ++ .{ .name = "drama3 nvfp4 d3_nvfp4_up", .threads = 128, .tg_memory = 0, .reserved = 128 }, ++ .{ .name = "drama3 nvfp4 d3_nvfp4_down", .threads = 128, .tg_memory = 0, .reserved = 128 }, ++ .{ .name = "drama3 nvfp4 d3_router_compact", .threads = 256, .tg_memory = 256, .reserved = 256 }, ++ .{ .name = "drama3 nvfp4 d3_nvfp4_combine", .threads = 256, .tg_memory = 128, .reserved = 256 }, ++ .{ .name = "drama3 guidance d3_cfg_mix", .threads = 256, .tg_memory = 0, .reserved = 256 }, + .{ .name = "drama3 sample d3_argmax", .threads = 1024, .tg_memory = 256, .reserved = 1024 }, + .{ .name = "drama3 sample d3_sample_main", .threads = 256, .tg_memory = 3632, .reserved = 256 }, + .{ .name = "drama3 sample d3_sample_code", .threads = 256, .tg_memory = 3632, .reserved = 256 }, +diff --git a/zig/src/core/engine_api.zig b/zig/src/core/engine_api.zig +index e1b5ece8..21a962d3 100644 +--- a/zig/src/core/engine_api.zig ++++ b/zig/src/core/engine_api.zig +@@ -264,6 +264,10 @@ pub const Open = struct { + context: ?i64 = null, + max_frames: ?u32 = null, + silence: speech_guard.Config = .{}, ++ /// Allocate an independent backbone state per speech lane; Metal Drama only. ++ cfg_enabled: bool = false, ++ /// Submit the two independent CFG sequences in one Metal command buffer, not vectorized GEMM. ++ cfg_batched: bool = false, + lanes: u32 = 8, + /// --parallel named a number: an engine whose memory fits fewer streams refuses instead of serving fewer. + lanes_fixed: bool = false, +diff --git a/zig/src/core/speech_api.zig b/zig/src/core/speech_api.zig +index d593239a..08ac16a7 100644 +--- a/zig/src/core/speech_api.zig ++++ b/zig/src/core/speech_api.zig +@@ -16,6 +16,9 @@ pub const SpeechRequest = struct { + stream: bool = false, + /// End of the invariant reference prompt, excluding request text; 0 disables reuse; backend saves no lookahead row. + reference_rows: u32 = 0, ++ /// Plain text frame rows for reference-conditioned semantic CFG; empty means no guidance. ++ unconditioned: []const [11]i32 = &.{}, ++ cfg_gamma: f32 = 1, + }; + + pub const SpeechInfo = struct { +diff --git a/zig/src/families/drama3/drama3_metal.zig b/zig/src/families/drama3/drama3_metal.zig +index 41b2e059..444f17d7 100644 +--- a/zig/src/families/drama3/drama3_metal.zig ++++ b/zig/src/families/drama3/drama3_metal.zig +@@ -16,6 +16,8 @@ pub const kernels = @import("metal/kernels.zig"); + pub const weights = @import("metal/weights.zig"); + pub const load_limit = @import("metal/load_limit.zig"); + pub const engine = @import("metal/engine.zig"); ++pub const guidance = @import("metal/guidance.zig"); ++pub const quant_check = @import("metal/quant_check.zig"); + pub const codec = @import("metal/codec.zig"); + pub const tts = @import("metal/tts.zig"); + +@@ -39,6 +41,8 @@ test { + _ = wav; + _ = kernels; + _ = load_limit; ++ _ = guidance; ++ _ = @import("metal/quant_format_test.zig"); + } + + pub const residency = @import("metal/residency.zig"); +diff --git a/zig/src/families/drama3/metal/engine.zig b/zig/src/families/drama3/metal/engine.zig +index 54f761c3..5ce06ebc 100644 +--- a/zig/src/families/drama3/metal/engine.zig ++++ b/zig/src/families/drama3/metal/engine.zig +@@ -108,7 +108,7 @@ pub const Engine = struct { + } + e.state_start = e.g.bufs.items.len; + errdefer e.g.releaseFrom(e.state_start); +- try e.allocState(); ++ try e.allocState(false); + e.state_end = e.g.bufs.items.len; + if (o.warmup) { + var rows: [4][W]i32 = @splat(@splat(0)); +@@ -131,13 +131,23 @@ pub const Engine = struct { + } + + pub fn fork(base: *Engine) !*Engine { ++ return base.forkState(false); ++ } ++ ++ /// CFG's independent backbone state, without codec tower caches, output frames or tower scratch. ++ pub fn forkBackbone(base: *Engine) !*Engine { ++ return base.forkState(true); ++ } ++ ++ fn forkState(base: *Engine, backbone_only: bool) !*Engine { + const e = try base.gpa.create(Engine); + errdefer base.gpa.destroy(e); + e.* = base.*; + e.owns_model = false; ++ if (backbone_only) e.o.max_rows = 1; + e.state_start = e.g.bufs.items.len; + errdefer e.g.releaseFrom(e.state_start); +- try e.allocState(); ++ try e.allocState(backbone_only); + e.state_end = e.g.bufs.items.len; + return e; + } +@@ -151,12 +161,12 @@ pub const Engine = struct { + return e.g.alloc(bytes); + } + +- fn allocState(e: *Engine) !void { ++ fn allocState(e: *Engine, backbone_only: bool) !void { + const R = e.o.max_rows; + const ctxn = e.o.max_ctx; + const kvb = ctxn * cfg.kv_heads * cfg.head_dim * 2; + for (0..cfg.layers) |li| e.kv_bb[li] = if (e.c.full_attention[li]) .{ try e.alloc(kvb), try e.alloc(kvb) } else .{ 0, 0 }; +- for (0..cfg.tower_layers) |ti| e.kv_tw[ti] = .{ try e.alloc(kvb), try e.alloc(kvb) }; ++ for (0..cfg.tower_layers) |ti| e.kv_tw[ti] = if (backbone_only) .{ 0, 0 } else .{ try e.alloc(kvb), try e.alloc(kvb) }; + var n_gdn: usize = 0; + for (e.c.full_attention) |f| n_gdn += @intFromBool(!f); + const conv_b = cfg.conv_channels * 3 * 2; +@@ -176,13 +186,15 @@ pub const Engine = struct { + } + inline for (.{ "pos", "pos_bb", "step", "count" }) |f| @field(e, f) = try e.alloc(8); + e.bb_input = try e.alloc(W * 4); +- inline for (.{ "t_xn", "t_tmp", "x", "xn", "tmp", "shout", "hfin", "mixed" }) |f| @field(e, f) = try e.alloc(R * H * 2); +- e.t_proj = try e.alloc(R * 2 * cfg.tower_ffn * 2); +- inline for (.{ "t_y", "t_q", "y", "q" }) |f| @field(e, f) = try e.alloc(R * 4096 * 2); ++ inline for (.{ "t_xn", "t_tmp", "mixed" }) |f| @field(e, f) = if (backbone_only) 0 else try e.alloc(R * H * 2); ++ inline for (.{ "x", "xn", "tmp", "shout", "hfin" }) |f| @field(e, f) = try e.alloc(R * H * 2); ++ e.t_proj = if (backbone_only) 0 else try e.alloc(R * 2 * cfg.tower_ffn * 2); ++ inline for (.{ "t_y", "t_q" }) |f| @field(e, f) = if (backbone_only) 0 else try e.alloc(R * 4096 * 2); ++ inline for (.{ "y", "q" }) |f| @field(e, f) = try e.alloc(R * 4096 * 2); + e.window = try e.alloc(16 * 4); + e.cur_input = try e.alloc(W * 4); +- e.frame = try e.alloc(W * 4); +- e.frames_out = try e.alloc(e.o.max_frames * W * 4); ++ e.frame = if (backbone_only) 0 else try e.alloc(W * 4); ++ e.frames_out = if (backbone_only) 0 else try e.alloc(e.o.max_frames * W * 4); + e.pframes = try e.alloc((ctxn + 1) * W * 4); + e.proj = try e.alloc(R * cfg.gdn_proj * 2); + e.conv_out = try e.alloc(R * cfg.conv_channels * 2); +@@ -199,12 +211,12 @@ pub const Engine = struct { + e.gsc = try e.alloc(R * 8 * (cfg.expert_width / 128) * 4); + e.counts = try e.alloc(cfg.experts * 4); + e.members = try e.alloc(cfg.experts * R * 8 * 4); +- e.th = .{ try e.alloc(R * H * 2), try e.alloc(R * H * 2) }; +- e.mix = try e.alloc(R * 2 * H * 2); +- e.act = try e.alloc(R * cfg.tower_ffn * 2); ++ e.th = if (backbone_only) .{ 0, 0 } else .{ try e.alloc(R * H * 2), try e.alloc(R * H * 2) }; ++ e.mix = if (backbone_only) 0 else try e.alloc(R * 2 * H * 2); ++ e.act = if (backbone_only) 0 else try e.alloc(R * cfg.tower_ffn * 2); + e.hlast = try e.alloc(H * 2); + e.cand_logits = try e.alloc(e.w.n_cand * 4); +- e.code_logits = try e.alloc(cfg.codebook_size * 4); ++ e.code_logits = if (backbone_only) 0 else try e.alloc(cfg.codebook_size * 4); + } + + pub fn flush(e: *Engine) !void { +@@ -260,8 +272,9 @@ pub const Engine = struct { + try e.put(e.pframes, std.mem.sliceAsBytes(rows)); + } + +- pub fn promptChunk(e: *Engine, n: usize, c0: usize, t: usize, probe: ?Probe) !void { +- if (t == 0 or t > e.o.max_rows or c0 >= n or t > n - c0) return error.BadPromptChunk; ++ /// Backbone-only prompt chunk for the unconditional CFG sequence; no codec tower work. ++ pub fn backboneChunk(e: *Engine, c0: usize, t: usize, probe: ?Probe) !void { ++ if (t == 0 or t > e.o.max_rows) return error.BadPromptChunk; + const g = e.g; + var nb: [64]u8 = undefined; + try e.setPos(c0); +@@ -273,6 +286,12 @@ pub const Engine = struct { + } + k.rmsnorm(g, e.x, H, e.w.final_norm, e.hfin, H, t, H, cfg.eps, true); + if (probe) |p| try p.f(p.ctx, "backbone_final_norm", e.hfin, t, H, false); ++ } ++ ++ pub fn promptChunk(e: *Engine, n: usize, c0: usize, t: usize, probe: ?Probe) !void { ++ if (t == 0 or t > e.o.max_rows or c0 >= n or t > n - c0) return error.BadPromptChunk; ++ try e.backboneChunk(c0, t, probe); ++ const g = e.g; + // the tower over this chunk's rows that have a next prompt row + const tt = @min(c0 + t, n - 1) - c0; + if (tt > 0) try e.tower(tt, c0, e.hfin, probe); +diff --git a/zig/src/families/drama3/metal/forward.zig b/zig/src/families/drama3/metal/forward.zig +index 5f79ccd4..72f1f4bb 100644 +--- a/zig/src/families/drama3/metal/forward.zig ++++ b/zig/src/families/drama3/metal/forward.zig +@@ -54,24 +54,35 @@ pub fn backboneLayer(e: *Engine, ln: Lane, li: usize, t: usize) void { + moe(e, L, t); + } + +-/// The MoE block: router, gated shared expert and the routed FP8 experts, added to the residual stream. ++/// The MoE block: original-ID router, gated shared expert and packed FP8/NVFP4 experts, added to the residual. + fn moe(e: *Engine, L: *const wts.Layer, t: usize) void { + const g = e.g; + const mp = cfg.moe_proj; + k.linear(g, L.moe_in, e.xn, H, H, mp, e.sp, mp, t, false); +- k.router(g, e.sp, mp, 2 * cfg.shared_width, e.picks, e.rw, t); ++ if (L.expert_count == cfg.experts) ++ k.router(g, e.sp, mp, 2 * cfg.shared_width, e.picks, e.rw, t) ++ else ++ k.routerCompact(g, e.sp, mp, 2 * cfg.shared_width, e.picks, e.rw, L.expert_slots, t); + k.swiglu(g, e.sp, mp, 0, cfg.shared_width, cfg.shared_width, e.sact, t); + k.linear(g, L.shared_down, e.sact, cfg.shared_width, cfg.shared_width, H, e.shout, H, t, false); + const pairs = t * cfg.top_k; + const prompt = t > 1; + if (prompt) k.plan(g, e.picks, pairs, e.counts, e.members, e.o.max_rows * cfg.top_k); +- const r: k.Route = .{ .table = L.experts, .picks = e.picks, .counts = e.counts, .members = e.members, .cap = e.o.max_rows * cfg.top_k, .pairs = pairs, .prompt = prompt }; +- // activations to FP8 per row and 128 inputs before each expert projection, as the reference serves them +- k.actQuant(g, e.xn, H, e.xq, H, e.xsc, t, H); +- k.experts(g, true, e.xq, e.xsc, r, H, cfg.expert_width, e.gu); +- k.actQuant(g, e.gu, cfg.expert_width, e.gq, cfg.expert_width, e.gsc, pairs, cfg.expert_width); +- k.experts(g, false, e.gq, e.gsc, r, cfg.expert_width, H, e.down); +- k.moeCombine(g, e.x, e.down, e.rw, e.shout, e.sp, mp, 2 * cfg.shared_width + cfg.experts, H, t); ++ const r: k.Route = .{ .table = L.experts, .picks = e.picks, .counts = e.counts, .members = e.members, .cap = e.o.max_rows * cfg.top_k, .pairs = pairs, .prompt = prompt, .experts = L.expert_count }; ++ if (L.nvfp4) { ++ k.expertsNvfp4(g, true, e.xn, r, H, cfg.expert_width, e.gu); ++ k.expertsNvfp4(g, false, e.gu, r, cfg.expert_width, H, e.down); ++ } else { ++ // The legacy FP8 path retains its activation quantization and reduction arithmetic. ++ k.actQuant(g, e.xn, H, e.xq, H, e.xsc, t, H); ++ k.experts(g, true, e.xq, e.xsc, r, H, cfg.expert_width, e.gu); ++ k.actQuant(g, e.gu, cfg.expert_width, e.gq, cfg.expert_width, e.gsc, pairs, cfg.expert_width); ++ k.experts(g, false, e.gq, e.gsc, r, cfg.expert_width, H, e.down); ++ } ++ if (L.nvfp4) ++ k.moeCombineNvfp4(g, e.x, e.down, e.rw, e.shout, e.xn, L.moe_in + (2 * cfg.shared_width + cfg.experts) * H * 2, H, t) ++ else ++ k.moeCombine(g, e.x, e.down, e.rw, e.shout, e.sp, mp, 2 * cfg.shared_width + cfg.experts, H, t); + } + + /// Tower layer ti over t rows: hidden hin -> hout; label frames lf (rows lfs ints apart), input frames inf. +@@ -96,14 +107,21 @@ fn probeName(buf: []u8, comptime f: []const u8, i: usize) []const u8 { + return std.fmt.bufPrint(buf, f, .{i}) catch unreachable; + } + ++pub const Guidance = struct { unconditioned: u64, gamma: f32 }; ++ + /// A frame step's dispatches: sample, the tower, record, then the backbone row at the next position. + pub fn enqueueFrame(e: *Engine, smp: Sampling, probe: ?Probe, force: ?Engine.Force) !void { ++ return enqueueFrameGuided(e, smp, probe, force, null); ++} ++ ++pub fn enqueueFrameGuided(e: *Engine, smp: Sampling, probe: ?Probe, force: ?Engine.Force, guidance: ?Guidance) !void { + const g = e.g; + const sem_begin = e.c.ids.sem_begin; + const greedy = smp.greedy(); + const sp = k.sampling(smp); + if (!greedy and (smp.top_k < 1 or smp.top_k > 64)) return error.TopKUnsupportedOnMetal; + k.linear(g, e.w.cand_head, e.hlast, H, H, e.w.n_cand, e.cand_logits, e.w.n_cand, 1, true); ++ if (guidance) |guide| @import("guidance.zig").mix(e, guide.unconditioned, guide.gamma); + if (probe) |p| try p.f(p.ctx, "tower_decode0.candidate_logits", e.cand_logits, 1, e.w.n_cand, true); + if (force) |fc| { + k.argmax(g, e.cand_logits, e.w.n_cand, e.w.cand_ids, fc.picks, sem_begin, cfg.codebook_size); +diff --git a/zig/src/families/drama3/metal/gpu.zig b/zig/src/families/drama3/metal/gpu.zig +index 294d845b..fefffdba 100644 +--- a/zig/src/families/drama3/metal/gpu.zig ++++ b/zig/src/families/drama3/metal/gpu.zig +@@ -22,6 +22,11 @@ pub const P = enum { + d3_plan, + d3_expert_up, + d3_expert_down, ++ d3_nvfp4_up, ++ d3_nvfp4_down, ++ d3_nvfp4_combine, ++ d3_router_compact, ++ d3_cfg_mix, + d3_argmax, + d3_sample_main, + d3_sample_code, +diff --git a/zig/src/families/drama3/metal/guidance.zig b/zig/src/families/drama3/metal/guidance.zig +new file mode 100644 +index 00000000..516ec30f +--- /dev/null ++++ b/zig/src/families/drama3/metal/guidance.zig +@@ -0,0 +1,152 @@ ++//! Semantic-only CFG with independent recurrent/KV state and one shared checkpoint weight set. ++const std = @import("std"); ++const cfg = @import("../config.zig"); ++const Engine = @import("engine.zig").Engine; ++const k = @import("kernels.zig"); ++const fwd = @import("forward.zig"); ++const H = cfg.hidden; ++ ++pub fn validGamma(gamma: f32) bool { ++ return std.math.isFinite(gamma) and gamma >= 1 and gamma <= 5; ++} ++ ++pub fn mix(e: *Engine, unconditional: u64, gamma: f32) void { ++ mixLogits(e.g, e.cand_logits, unconditional, e.w.n_cand, gamma); ++} ++ ++fn mixLogits(g: *@import("gpu.zig").Gpu, conditional: u64, unconditional: u64, count: usize, gamma: f32) void { ++ const Args = extern struct { count: u32, gamma: f32 }; ++ g.go(.d3_cfg_mix, .{ (count + 255) / 256, 1, 1 }, 256, &.{ conditional, unconditional }, Args{ .count = @intCast(count), .gamma = gamma }); ++} ++ ++/// Host lifecycle is separate from GPU storage; no old reference may activate a later plain request. ++pub const State = struct { ++ active: bool = false, ++ gamma: f32 = 1, ++ ++ pub fn begin(s: *State, rows: usize, gamma: f32) !void { ++ s.clear(); ++ if (!validGamma(gamma)) return error.BadDramaGuidance; ++ if (rows == 0) return; ++ s.* = .{ .active = true, .gamma = gamma }; ++ } ++ ++ pub fn clear(s: *State) void { ++ s.* = .{}; ++ } ++}; ++ ++pub const Guide = struct { ++ e: *Engine, ++ batched: bool, ++ state: State = .{}, ++ ++ pub fn create(base: *Engine, batched: bool) !Guide { ++ return .{ .e = try base.forkBackbone(), .batched = batched }; ++ } ++ ++ pub fn destroy(g: *Guide) void { ++ g.e.destroy(); ++ } ++ ++ pub fn begin(g: *Guide, rows: []const [cfg.frame_width]i32, gamma: f32) !void { ++ try g.state.begin(rows.len, gamma); ++ errdefer g.state.clear(); ++ // A guide is never restored from the conditional reference cache. ++ if (g.state.active) try g.e.startPrompt(rows); ++ } ++ ++ /// One prompt row per scheduler round keeps cancellation and other lanes responsive. ++ pub fn promptStep(g: *Guide, rows: []const [cfg.frame_width]i32) !bool { ++ if (!g.state.active or g.e.rows_done == rows.len) return false; ++ try g.e.backboneChunk(g.e.rows_done, 1, null); ++ g.e.rows_done += 1; ++ if (g.e.rows_done == rows.len) { ++ g.e.g.copy(g.e.hfin, g.e.hlast, H * 2); ++ try g.e.finishPrompt(rows); ++ } ++ return true; ++ } ++ ++ pub fn clear(g: *Guide) void { ++ g.state.clear(); ++ } ++ /// Paired means one command-buffer submission, not a two-row GEMM or simultaneous GPU execution. ++ pub fn frame(g: *Guide, conditional: *Engine, sampling: @import("../common.zig").Sampling) !void { ++ if (!g.state.active) return conditional.frameStep(sampling, true, null); ++ const u = g.e; ++ k.linear(u.g, u.w.cand_head, u.hlast, H, H, u.w.n_cand, u.cand_logits, u.w.n_cand, 1, true); ++ if (!g.batched) try u.g.commit(); ++ try fwd.enqueueFrameGuided(conditional, sampling, null, null, .{ .unconditioned = u.cand_logits, .gamma = g.state.gamma }); ++ if (!g.batched) try u.g.commit(); ++ // Feed the actual sampled conditional frame to the unconditional backbone, never its codec tower. ++ k.nextInput(u.g, conditional.frame, u.bb_input, u.pos, u.pos_bb); ++ k.embedBackbone(u.g, u.bb_input, u.w.embed, u.w.cb_emb, u.x, 1, H, u.c.ids.sem_begin, u.c.ids.sem_end, cfg.codebook_size); ++ for (0..cfg.layers) |li| fwd.backboneLayer(u, fwd.laneBB(u, u.pos_bb), li, 1); ++ k.rmsnorm(u.g, u.x, H, u.w.final_norm, u.hlast, H, 1, H, cfg.eps, true); ++ k.copyI32(u.g, u.pos_bb, u.pos, 1); ++ try u.g.commit(); ++ } ++}; ++ ++test "CFG rejects nonfinite and out-of-range gamma before state activation" { ++ for ([_]f32{ 0, 0.99, 5.01, std.math.nan(f32), std.math.inf(f32) }) |g| ++ try std.testing.expect(!validGamma(g)); ++} ++ ++test "CFG transitions clear old voice on plain request cancellation and invalid begin" { ++ var state: State = .{}; ++ try state.begin(12, 1.5); ++ try std.testing.expect(state.active); ++ try state.begin(0, 1.5); ++ try std.testing.expect(!state.active and state.gamma == 1); ++ try state.begin(8, 2); ++ state.clear(); ++ try std.testing.expect(!state.active); ++ try state.begin(9, 3); ++ try std.testing.expectError(error.BadDramaGuidance, state.begin(9, std.math.nan(f32))); ++ try std.testing.expect(!state.active and state.gamma == 1); ++ try state.begin(7, 1); ++ try std.testing.expect(state.active and state.gamma == 1); ++} ++ ++/// Synthetic device qualification called only by the explicit Metal quant-check executable. ++pub fn check(g: *@import("gpu.zig").Gpu) !void { ++ const first = g.bufs.items.len; ++ defer g.releaseFrom(first); ++ const n = cfg.codebook_size + 1; ++ // The trailing sentinels model codec logits, which must never be touched by semantic CFG. ++ var conditional: [n + 8]f32 = undefined; ++ var unconditional: [n + 8]f32 = undefined; ++ var actual: [n + 8]f32 = undefined; ++ for (&conditional, &unconditional, 0..) |*c, *u, i| { ++ c.* = @as(f32, @floatFromInt(@as(i32, @intCast(i % 97)) - 48)) / 8; ++ u.* = @as(f32, @floatFromInt(@as(i32, @intCast(i % 73)) - 36)) / 16; ++ } ++ conditional[0] = -0.0; ++ const cbuf = try g.alloc(@sizeOf(@TypeOf(conditional))); ++ const ubuf = try g.alloc(@sizeOf(@TypeOf(unconditional))); ++ for ([_]f32{ 1, 1.5, 5 }) |gamma| { ++ // Gamma one must preserve bits even if an unused unconditional value is NaN. ++ unconditional[0] = if (gamma == 1) std.math.nan(f32) else -2; ++ try g.put(cbuf, std.mem.asBytes(&conditional)); ++ try g.put(ubuf, std.mem.asBytes(&unconditional)); ++ mixLogits(g, cbuf, ubuf, n, gamma); ++ try g.get(cbuf, std.mem.asBytes(&actual)); ++ for (actual, conditional, unconditional, 0..) |got, c, u, i| { ++ if (gamma == 1 or i >= n) { ++ if (@as(u32, @bitCast(got)) != @as(u32, @bitCast(c))) return error.GuidanceIdentityMismatch; ++ } else { ++ const expected = @as(f64, c) + (@as(f64, gamma) - 1) * (@as(f64, c) - @as(f64, u)); ++ if (!std.math.isFinite(got) or @abs(@as(f64, got) - expected) > 1e-6) ++ return error.GuidanceFormulaMismatch; ++ } ++ } ++ } ++ for ([_]f32{ std.math.nan(f32), std.math.inf(f32), -std.math.inf(f32), 0, 5.01 }) |bad| { ++ var state: State = .{}; ++ if (state.begin(2, bad)) |_| return error.GuidanceAcceptedInvalidGamma else |err| { ++ if (err != error.BadDramaGuidance or state.active) return error.GuidanceAcceptedInvalidGamma; ++ } ++ } ++} +diff --git a/zig/src/families/drama3/metal/kernels.zig b/zig/src/families/drama3/metal/kernels.zig +index cb6311cb..18957311 100644 +--- a/zig/src/families/drama3/metal/kernels.zig ++++ b/zig/src/families/drama3/metal/kernels.zig +@@ -50,6 +50,10 @@ pub fn moeCombine(g: *Gpu, r: u64, down: u64, w: u64, shared: u64, sp: u64, ss: + g.go(.d3_moe_combine, .{ t, 1, 1 }, 256, &.{ r, down, w, shared, sp }, A{ .ss = u(ss), .sg = u(sg_off), .D = u(dim) }); + } + ++pub fn moeCombineNvfp4(g: *Gpu, residual: u64, down: u64, weights: u64, shared: u64, x: u64, gate: u64, dim: usize, rows: usize) void { ++ g.go(.d3_nvfp4_combine, .{ rows, 1, 1 }, 256, &.{ residual, down, weights, shared, x, gate }, u(dim)); ++} ++ + pub const Mix = struct { lf: u64, lfs: usize, code_col: usize, inf: u64, with_shared: bool, table: u64, cb_emb: u64, h: u64, enorm: u64, hnorm: u64, out: u64 }; + + pub fn towerMix(g: *Gpu, m: Mix, dim: usize, sem_begin: i32, sem_end: i32, cb: i32, eps: f32, t: usize) void { +@@ -89,6 +93,11 @@ pub fn router(g: *Gpu, sp: u64, ss: usize, r_off: usize, picks: u64, weights: u6 + g.go(.d3_router, .{ t, 1, 1 }, 256, &.{ sp, picks, weights }, A{ .ss = u(ss), .off = u(r_off) }); + } + ++pub fn routerCompact(g: *Gpu, sp: u64, ss: usize, r_off: usize, picks: u64, weights: u64, slots: u64, t: usize) void { ++ const A = extern struct { ss: u32, off: u32 }; ++ g.go(.d3_router_compact, .{ t, 1, 1 }, 256, &.{ sp, picks, weights, slots }, A{ .ss = u(ss), .off = u(r_off) }); ++} ++ + /// x rows to FP8 values (kept in bf16) and their scales [rows, kdim / 128]. + pub fn actQuant(g: *Gpu, x: u64, xs: usize, q: u64, qs: usize, sc: u64, rows: usize, kdim: usize) void { + const A = extern struct { xs: u32, qs: u32, groups: u32 }; +@@ -100,16 +109,23 @@ pub fn plan(g: *Gpu, picks: u64, pairs: usize, counts: u64, members: u64, cap: u + g.go(.d3_plan, .{ 1, 1, 1 }, 1024, &.{ picks, counts, members }, A{ .pairs = u(pairs), .cap = u(cap) }); + } + +-pub const Route = struct { table: u64, picks: u64, counts: u64, members: u64, cap: usize, pairs: usize, prompt: bool }; ++pub const Route = struct { table: u64, picks: u64, counts: u64, members: u64, cap: usize, pairs: usize, prompt: bool, experts: usize = 256 }; + + /// Routed experts: `up` maps token rows (x) to SwiGLU pair rows, else pair rows to fp32 down rows. + pub fn experts(g: *Gpu, up: bool, x: u64, xsc: u64, r: Route, kdim: usize, n: usize, out: u64) void { + const A = extern struct { K: u32, N: u32, prompt: u32, cap: u32, slots: u32 }; + const a: A = .{ .K = u(kdim), .N = u(n), .prompt = @intFromBool(r.prompt), .cap = u(r.cap), .slots = 8 }; +- const gy: usize = if (r.prompt) 256 else r.pairs; ++ const gy: usize = if (r.prompt) r.experts else r.pairs; + g.go(if (up) .d3_expert_up else .d3_expert_down, .{ n / @as(usize, if (up) 8 else 16), gy, 1 }, 128, &.{ x, xsc, r.table, r.picks, r.counts, r.members, out }, a); + } + ++/// NVFP4 reads BF16 activation rows directly and writes BF16 SwiGLU / FP32 down; no activation quantization. ++pub fn expertsNvfp4(g: *Gpu, up: bool, x: u64, r: Route, kdim: usize, n: usize, out: u64) void { ++ const A = extern struct { K: u32, N: u32, prompt: u32, cap: u32, slots: u32 }; ++ const a: A = .{ .K = u(kdim), .N = u(n), .prompt = @intFromBool(r.prompt), .cap = u(r.cap), .slots = 8 }; ++ g.go(if (up) .d3_nvfp4_up else .d3_nvfp4_down, .{ n / 4, if (r.prompt) r.experts else r.pairs, 1 }, 128, &.{ x, r.table, r.picks, r.counts, r.members, out }, a); ++} ++ + /// Greedy: cand_ids non-zero picks a main token (frame[0..2]), zero a code (frame[0]). + pub fn argmax(g: *Gpu, logits: u64, count: usize, cand_ids: u64, frame: u64, sem_begin: i32, cb: i32) void { + const A = extern struct { count: u32, ids: u32, sb: i32, cb: i32 }; +diff --git a/zig/src/families/drama3/metal/load_limit.zig b/zig/src/families/drama3/metal/load_limit.zig +index 0e024444..957dd109 100644 +--- a/zig/src/families/drama3/metal/load_limit.zig ++++ b/zig/src/families/drama3/metal/load_limit.zig +@@ -60,6 +60,11 @@ pub fn cacheBytes(c: cfg.Config, context: u32) u64 { + + /// The bytes a load needs: the checkpoint's tensors (read in place, the bf16 ones copied too) and the caches. + pub fn planBytes(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, context: u32) !u64 { ++ return planBytesStates(gpa, io, dir, context, 1); ++} ++ ++pub fn planBytesStates(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, context: u32, states: u32) !u64 { ++ if (states == 0) return error.BadStateCount; + const c = try cfg.Config.load(gpa, io, dir); + var sh = try shards.Shards.open(gpa, io, dir); + defer sh.close(io); +@@ -67,9 +72,9 @@ pub fn planBytes(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, context: u + for (sh.files) |f| for (f.names.values()) |e| { + const bytes: u64 = e.end - e.begin; + resident += bytes; +- if (e.dtype == .bf16) resident += bytes; // copied to the device beside the mapped pages ++ if (e.dtype == .bf16) resident += bytes; + }; +- return resident + cacheBytes(c, context); ++ return resident + cacheBytes(c, context) * states; + } + + /// A plan past the limit is refused before a weight is read. +@@ -79,10 +84,15 @@ pub fn admitUnder(plan: u64, limit: Limit) error{OverMemoryLimit}!void { + + /// This Mac's ceiling and the load's plan: a load that does not fit is refused before a weight is read. + pub fn admit(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, context: u32) !void { ++ return admitStates(gpa, io, dir, context, 1); ++} ++ ++/// Independent guidance/serving states share weights, but each owns its caches and row scratch. ++pub fn admitStates(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, context: u32, states: u32) !void { + const device = try mtl.Device.init(); + defer device.deinit(); + const limit = try loadLimit(device); +- const plan = try planBytes(gpa, io, dir, context); ++ const plan = try planBytesStates(gpa, io, dir, context, states); + if (plan > limit.bytes) { + std.log.err("drama3: {d:.1} GB of weights and caches pass this Mac's {d:.1} GB load limit ({s}); free memory, or set DRAMA_LOAD_LIMIT_GB (past 70% of RAM a load risks a watchdog panic)", .{ toGb(plan), toGb(limit.bytes), limit.source() }); + return error.OverMemoryLimit; +diff --git a/zig/src/families/drama3/metal/quant_check.zig b/zig/src/families/drama3/metal/quant_check.zig +new file mode 100644 +index 00000000..36d4f264 +--- /dev/null ++++ b/zig/src/families/drama3/metal/quant_check.zig +@@ -0,0 +1,232 @@ ++//! Deterministic device gate: packed projections, SwiGLU, pruning and mixed FP8/NVFP4 vs independent FP64. ++const std = @import("std"); ++const Gpu = @import("gpu.zig").Gpu; ++const k = @import("kernels.zig"); ++const q = @import("quant_format.zig"); ++const ref = @import("quant_reference.zig"); ++ ++const H = 2048; ++const I = 512; ++const E = 8; ++const max_rows = 3; ++const pairs = max_rows * E; ++ ++const Score = struct { ++ count: usize = 0, ++ worst: f64 = 0, ++ worst_column: usize = 0, ++ squared: f64 = 0, ++ fn add(s: *Score, got: f64, want: f64, bound: f64, col: usize) !void { ++ if (!std.math.isFinite(got) or !std.math.isFinite(want)) return error.NonfiniteQuantResult; ++ const delta = @abs(got - want); ++ s.count += 1; ++ s.squared += delta * delta; ++ if (delta > s.worst) { ++ s.worst = delta; ++ s.worst_column = col; ++ } ++ if (delta > bound) { ++ std.debug.print("FAIL quant column={d} got={d} FP64={d} abs={d} bound={d}\n", .{ col, got, want, delta, bound }); ++ return error.QuantReferenceMismatch; ++ } ++ } ++ fn report(s: Score, label: []const u8) void { ++ std.debug.print("SUMMARY {s} PASS count={d} max_abs={d} worst_column={d} rms={d}\n", .{ label, s.count, s.worst, s.worst_column, @sqrt(s.squared / @as(f64, @floatFromInt(s.count))) }); ++ } ++}; ++ ++fn put(g: *Gpu, bytes: []const u8, indirect: bool) !u64 { ++ const a = try g.alloc(bytes.len); ++ @memcpy(g.host(a, bytes.len), bytes); ++ if (indirect) try g.indirect.append(g.gpa, g.ref(a).buf); ++ return a; ++} ++ ++fn bfRows(g: *Gpu, a: u64, n: usize) []align(1) u16 { ++ return std.mem.bytesAsSlice(u16, @as([]align(1) u8, g.host(a, n * 2))); ++} ++ ++const Projection = struct { ++ w: u64, ++ s: u64, ++ global: u64, ++ n: usize, ++ dim: usize, ++ fn init(g: *Gpu, n: usize, dim: usize, seed: usize) !Projection { ++ const w = try g.alloc(n * dim / 2); ++ const s = try g.alloc(n * dim / 16); ++ for (g.host(w, n * dim / 2), 0..) |*b, j| b.* = @truncate((j * 37 + seed * 19) % 256); ++ for (g.host(s, n * dim / 16), 0..) |*b, j| b.* = @intCast((j + seed * 11) % 127); ++ const global: f32 = @as(f32, @floatFromInt(seed + 1)) / 1024; ++ const gp = try put(g, std.mem.asBytes(&global), true); ++ try g.indirect.append(g.gpa, g.ref(w).buf); ++ try g.indirect.append(g.gpa, g.ref(s).buf); ++ return .{ .w = w, .s = s, .global = gp, .n = n, .dim = dim }; ++ } ++ fn scalar(p: Projection, g: *Gpu, row: usize, x: []align(1) const u16) f64 { ++ const global: f32 = @bitCast(std.mem.readInt(u32, g.host(p.global, 4)[0..4], .little)); ++ return ref.projection(g.host(p.w, p.n * p.dim / 2), g.host(p.s, p.n * p.dim / 16), global, row, p.dim, x); ++ } ++}; ++ ++/// Run only in an exclusive GPU window. The caller opens the newly built drama3 metallib, not a model checkpoint. ++pub fn run(g: *Gpu) !void { ++ try g.flush(); ++ const first = g.bufs.items.len; ++ const indirect_first = g.indirect.items.len; ++ defer { ++ g.flush() catch {}; ++ g.indirect.shrinkRetainingCapacity(indirect_first); ++ g.releaseFrom(first); ++ } ++ const gate = try Projection.init(g, I, H, 0); ++ const up = try Projection.init(g, I, H, 1); ++ const down = try Projection.init(g, H, I, 2); ++ var pointers: [E * 9]u64 = undefined; ++ for (0..E) |e| for ([_]Projection{ gate, up, down }, 0..) |p, j| { ++ pointers[e * 9 + j * 3] = g.gpuAddress(p.w); ++ pointers[e * 9 + j * 3 + 1] = g.gpuAddress(p.s); ++ pointers[e * 9 + j * 3 + 2] = g.gpuAddress(p.global); ++ }; ++ const table = try put(g, std.mem.sliceAsBytes(&pointers), false); ++ const x = try g.alloc(max_rows * H * 2); ++ for (bfRows(g, x, max_rows * H), 0..) |*b, j| b.* = q.bf16(@as(f32, @floatFromInt(@as(i32, @intCast((j * 13) % 29)) - 14)) / 512); ++ const logits = try g.alloc(max_rows * 256 * 2); ++ const original = [_]usize{ 1, 7, 19, 32, 65, 127, 200, 255 }; ++ var slot_map: [256]i32 = @splat(-1); ++ for (original, 0..) |id, slot| slot_map[id] = @intCast(slot); ++ const slots = try put(g, std.mem.sliceAsBytes(&slot_map), false); ++ // Excluded experts deliberately have the largest logits; retained ones tie, proving original-ID tie order. ++ for (bfRows(g, logits, max_rows * 256), 0..) |*b, j| b.* = q.bf16(if (slot_map[j % 256] >= 0) -2 else 100); ++ const picks = try g.alloc(pairs * 4); ++ const rw = try g.alloc(pairs * 4); ++ const counts = try g.alloc(256 * 4); ++ const members = try g.alloc(256 * pairs * 4); ++ const gu = try g.alloc(pairs * I * 2); ++ const output = try g.alloc(pairs * H * 4); ++ var up_score: Score = .{}; ++ var down_score: Score = .{}; ++ for ([_]usize{ 1, max_rows }) |rows| { ++ k.routerCompact(g, logits, 256, 0, picks, rw, slots, rows); ++ if (rows > 1) k.plan(g, picks, rows * E, counts, members, pairs); ++ const route: k.Route = .{ .table = table, .picks = picks, .counts = counts, .members = members, .cap = pairs, .pairs = rows * E, .prompt = rows > 1, .experts = E }; ++ k.expertsNvfp4(g, true, x, route, H, I, gu); ++ k.expertsNvfp4(g, false, gu, route, I, H, output); ++ try g.flush(); ++ for (0..rows * E) |pair| { ++ const picked = std.mem.readInt(i32, g.host(picks + pair * 4, 4)[0..4], .little); ++ if (picked != pair % E) return error.PrunedRoutingMismatch; ++ const weight: f32 = @bitCast(std.mem.readInt(u32, g.host(rw + pair * 4, 4)[0..4], .little)); ++ if (weight != 0.125) return error.PrunedRoutingMismatch; ++ const xr = bfRows(g, x + (pair / E) * H * 2, H); ++ const gr = bfRows(g, gu + pair * I * 2, I); ++ for (0..I) |col| { ++ const want = ref.swiglu(gate.scalar(g, col, xr), up.scalar(g, col, xr)); ++ try up_score.add(ref.fromBf16(gr[col]), want, 0.000002 + @abs(want) / 64, col); ++ } ++ for (0..H) |col| { ++ const want = down.scalar(g, col, gr); ++ const got: f32 = @bitCast(std.mem.readInt(u32, g.host(output + (pair * H + col) * 4, 4)[0..4], .little)); ++ try down_score.add(got, want, 0.0002 + @abs(want) * 0.00002, col); ++ } ++ } ++ try combine(g, output, rw, x, rows); ++ try rawProjection(g, route, gate, x, output, rows, "drama3 NVFP4 gate FP32"); ++ try rawProjection(g, route, up, x, output, rows, "drama3 NVFP4 up FP32"); ++ try fp8Layer(g, route, x, gu, output, rows); ++ } ++ up_score.report("drama3 NVFP4 gate/up SwiGLU decode+prompt"); ++ down_score.report("drama3 NVFP4 down FP32 decode+prompt"); ++} ++ ++fn combine(g: *Gpu, down: u64, weights: u64, x: u64, rows: usize) !void { ++ const residual = try g.alloc(rows * H * 2); ++ const shared = try g.alloc(rows * H * 2); ++ const wg = try g.alloc(H * 2); ++ for (bfRows(g, residual, rows * H), 0..) |*v, i| v.* = q.bf16(@as(f32, @floatFromInt(i % 13)) / 32); ++ for (bfRows(g, shared, rows * H), 0..) |*v, i| v.* = q.bf16(@as(f32, @floatFromInt(@as(i32, @intCast(i % 17)) - 8)) / 16); ++ for (bfRows(g, wg, H), 0..) |*v, i| v.* = q.bf16(@as(f32, @floatFromInt(@as(i32, @intCast(i % 11)) - 5)) / 64); ++ k.moeCombineNvfp4(g, residual, down, weights, shared, x, wg, H, rows); ++ try g.flush(); ++ var score: Score = .{}; ++ for (0..rows) |row| { ++ var dot: f64 = 0; ++ for (0..H) |j| dot += ref.fromBf16(bfRows(g, x, rows * H)[row * H + j]) * ref.fromBf16(bfRows(g, wg, H)[j]); ++ const gate = 1 / (1 + @exp(-dot)); ++ for (0..H) |col| { ++ var sum: f64 = 0; ++ for (0..E) |slot| { ++ const weight: f32 = @bitCast(std.mem.readInt(u32, g.host(weights + (row * E + slot) * 4, 4)[0..4], .little)); ++ const d: f32 = @bitCast(std.mem.readInt(u32, g.host(down + ((row * E + slot) * H + col) * 4, 4)[0..4], .little)); ++ sum += ref.roundBf16(@as(f32, weight * d)); ++ } ++ const sh = ref.fromBf16(bfRows(g, shared, rows * H)[row * H + col]); ++ const old = ref.roundBf16(@as(f64, @floatFromInt((row * H + col) % 13)) / 32); ++ const want = ref.roundBf16(old + ref.roundBf16(gate * sh + ref.roundBf16(sum))); ++ const got = ref.fromBf16(bfRows(g, residual, rows * H)[row * H + col]); ++ try score.add(got, want, 0.00001 + @abs(want) / 128, col); ++ } ++ } ++ score.report("drama3 NVFP4 pair rounding shared gate residual"); ++} ++ ++fn rawProjection(g: *Gpu, original_route: k.Route, p: Projection, x: u64, output: u64, rows: usize, label: []const u8) !void { ++ var ptrs: [E * 9]u64 = @splat(0); ++ for (0..E) |e| { ++ ptrs[e * 9 + 6] = g.gpuAddress(p.w); ++ ptrs[e * 9 + 7] = g.gpuAddress(p.s); ++ ptrs[e * 9 + 8] = g.gpuAddress(p.global); ++ } ++ var route = original_route; ++ route.table = try put(g, std.mem.sliceAsBytes(&ptrs), false); ++ const repeated = try g.alloc(rows * E * H * 2); ++ for (0..rows * E) |pair| ++ @memcpy(g.host(repeated + pair * H * 2, H * 2), g.host(x + (pair / E) * H * 2, H * 2)); ++ k.expertsNvfp4(g, false, repeated, route, H, I, output); ++ try g.flush(); ++ var score: Score = .{}; ++ for (0..rows * E) |pair| for (0..I) |col| { ++ const want = p.scalar(g, col, bfRows(g, repeated + pair * H * 2, H)); ++ const got: f32 = @bitCast(std.mem.readInt(u32, g.host(output + (pair * I + col) * 4, 4)[0..4], .little)); ++ try score.add(got, want, 0.00002 + @abs(want) * 0.00002, col); ++ }; ++ score.report(label); ++} ++ ++fn fp8Layer(g: *Gpu, original_route: k.Route, x: u64, gu: u64, output: u64, rows: usize) !void { ++ const w = try g.alloc(H * I); ++ for (g.host(w, H * I), 0..) |*b, j| b.* = @intCast(0x20 + j % 32); ++ const scales = try g.alloc((H / 128) * (I / 128) * 4); ++ const fs = std.mem.bytesAsSlice(f32, @as([]align(1) u8, g.host(scales, (H / 128) * (I / 128) * 4))); ++ @memset(fs, 0.0078125); ++ try g.indirect.append(g.gpa, g.ref(w).buf); ++ try g.indirect.append(g.gpa, g.ref(scales).buf); ++ var ptrs: [E * 6]u64 = undefined; ++ for (0..E * 3) |i| { ++ ptrs[i * 2] = g.gpuAddress(w); ++ ptrs[i * 2 + 1] = g.gpuAddress(scales); ++ } ++ var route = original_route; ++ route.table = try put(g, std.mem.sliceAsBytes(&ptrs), false); ++ const xq = try g.alloc(rows * H * 2); ++ const xs = try g.alloc(rows * (H / 128) * 4); ++ const gq = try g.alloc(rows * E * I * 2); ++ const gs = try g.alloc(rows * E * (I / 128) * 4); ++ k.actQuant(g, x, H, xq, H, xs, rows, H); ++ k.experts(g, true, xq, xs, route, H, I, gu); ++ k.actQuant(g, gu, I, gq, I, gs, rows * E, I); ++ k.experts(g, false, gq, gs, route, I, H, output); ++ try g.flush(); ++ var score: Score = .{}; ++ for (0..rows * E) |pair| for (0..H) |col| { ++ var want: f64 = 0; ++ const act = bfRows(g, gq + pair * I * 2, I); ++ for (0..I) |j| { ++ const scale: f32 = @bitCast(std.mem.readInt(u32, g.host(gs + (pair * (I / 128) + j / 128) * 4, 4)[0..4], .little)); ++ want += ref.fp8(g.host(w, H * I)[col * I + j]) * ref.fromBf16(act[j]) * scale * 0.0078125; ++ } ++ const got: f32 = @bitCast(std.mem.readInt(u32, g.host(output + (pair * H + col) * 4, 4)[0..4], .little)); ++ try score.add(got, want, 0.0001 + @abs(want) * 0.00002, col); ++ }; ++ score.report("drama3 mixed-layer legacy FP8 down"); ++} +diff --git a/zig/src/families/drama3/metal/quant_format.zig b/zig/src/families/drama3/metal/quant_format.zig +new file mode 100644 +index 00000000..a3e4959b +--- /dev/null ++++ b/zig/src/families/drama3/metal/quant_format.zig +@@ -0,0 +1,73 @@ ++//! Packed expert metadata: checkpoint tensor headers, not config quantization labels, select the format. ++const std = @import("std"); ++const core = @import("core"); ++const cfg = @import("../config.zig"); ++const Tensor = core.safetensors.Tensor; ++ ++pub const Held = struct { ++ counts: [cfg.layers]u16 = @splat(cfg.experts), ++ ids: [cfg.layers][cfg.experts]u8 = @splat(identity), ++ pub const identity = blk: { ++ var ids: [cfg.experts]u8 = undefined; ++ for (&ids, 0..) |*id, i| id.* = @intCast(i); ++ break :blk ids; ++ }; ++ pub fn slots(h: *const Held, layer: usize) [cfg.experts]i32 { ++ var out: [cfg.experts]i32 = @splat(-1); ++ for (h.ids[layer][0..h.counts[layer]], 0..) |id, slot| out[id] = @intCast(slot); ++ return out; ++ } ++}; ++ ++pub fn parseHeld(value: std.json.Value) !Held { ++ if (value != .object or value.object.count() != cfg.layers) return error.BadExpertMap; ++ var h: Held = .{}; ++ for (0..cfg.layers) |li| { ++ var key: [8]u8 = undefined; ++ const ids = value.object.get(try std.fmt.bufPrint(&key, "{d}", .{li})) orelse return error.BadExpertMap; ++ if (ids != .array or ids.array.items.len < cfg.top_k or ids.array.items.len > cfg.experts) return error.BadExpertMap; ++ for (ids.array.items, 0..) |id, s| { ++ if (id != .integer or id.integer < 0 or id.integer >= cfg.experts) return error.BadExpertMap; ++ if (s > 0 and id.integer <= h.ids[li][s - 1]) return error.BadExpertMap; ++ h.ids[li][s] = @intCast(id.integer); ++ } ++ h.counts[li] = @intCast(ids.array.items.len); ++ } ++ return h; ++} ++ ++pub fn heldMetadata(metadata: ?std.json.Value) !Held { ++ const meta = metadata orelse return .{}; ++ if (meta != .object) return error.BadExpertMap; ++ if (meta.object.get("drama_head_rows") != null) return error.UnsupportedCompactHead; ++ return if (meta.object.get("drama_experts")) |v| try parseHeld(v) else .{}; ++} ++ ++pub fn checkNames(h: *const Held, names: anytype) !void { ++ var it = names; ++ const pre = "model.language_model.layers."; ++ while (it.next()) |key| { ++ const name = key.*; ++ if (!std.mem.startsWith(u8, name, pre)) continue; ++ var parts = std.mem.splitScalar(u8, name[pre.len..], '.'); ++ const li = parts.next() orelse continue; ++ if (!std.mem.eql(u8, parts.next() orelse continue, "mlp")) continue; ++ if (!std.mem.eql(u8, parts.next() orelse continue, "experts")) continue; ++ const l = std.fmt.parseInt(usize, li, 10) catch return error.BadExpertMap; ++ const e = std.fmt.parseInt(u8, parts.next() orelse return error.BadExpertMap, 10) catch return error.BadExpertMap; ++ if (l >= cfg.layers or std.mem.indexOfScalar(u8, h.ids[l][0..h.counts[l]], e) == null) return error.UnlistedExpert; ++ } ++} ++ ++pub fn nvfp4(w: Tensor, s: Tensor, g: Tensor, n: usize, k: usize) !void { ++ if (k % 16 != 0 or !w.is(.u8, &.{ n, k / 2 }) or !s.is(.f8_e4m3, &.{ n, k / 16 }) or !g.is(.f32, &.{})) return error.BadNvfp4Shape; ++ if (w.bytes.len != n * k / 2 or s.bytes.len != n * k / 16 or g.bytes.len != 4) return error.BadNvfp4Shape; ++ for (s.bytes) |b| if (b == 0x7f or b > 0x80) return error.BadNvfp4Scale; ++ const global: f32 = @bitCast(std.mem.readInt(u32, g.bytes[0..4], .little)); ++ if (!(global > 0 and std.math.isFinite(global))) return error.BadScale; ++} ++ ++pub fn bf16(x: f32) u16 { ++ const bits: u32 = @bitCast(x); ++ return @truncate((bits +% (0x7fff + ((bits >> 16) & 1))) >> 16); ++} +diff --git a/zig/src/families/drama3/metal/quant_format_test.zig b/zig/src/families/drama3/metal/quant_format_test.zig +new file mode 100644 +index 00000000..747124a5 +--- /dev/null ++++ b/zig/src/families/drama3/metal/quant_format_test.zig +@@ -0,0 +1,90 @@ ++const std = @import("std"); ++const st = @import("core").safetensors; ++const q = @import("quant_format.zig"); ++const ref = @import("quant_reference.zig"); ++const cfg = @import("../config.zig"); ++ ++fn tensor(dtype: st.DType, shape: []const usize, bytes: []const u8) st.Tensor { ++ var t: st.Tensor = .{ .dtype = dtype, .rank = @intCast(shape.len), .shape = @splat(1), .bytes = bytes }; ++ @memcpy(t.shape[0..shape.len], shape); ++ return t; ++} ++ ++test "NVFP4 rejects shape dtype negative nan scales and invalid globals" { ++ const codes: [32]u8 = @splat(0x98); ++ var scales: [4]u8 = @splat(0x38); ++ var global: f32 = 1; ++ const w = tensor(.u8, &.{ 1, 32 }, &codes); ++ const s = tensor(.f8_e4m3, &.{ 1, 4 }, &scales); ++ const g = tensor(.f32, &.{}, std.mem.asBytes(&global)); ++ try q.nvfp4(w, s, g, 1, 64); ++ for ([_]u8{ 0x7f, 0x81, 0xff }) |bad| { ++ scales[0] = bad; ++ try std.testing.expectError(error.BadNvfp4Scale, q.nvfp4(w, s, g, 1, 64)); ++ } ++ for ([_]u8{ 0, 0x80 }) |zero| { ++ scales[0] = zero; ++ try q.nvfp4(w, s, g, 1, 64); ++ } ++ for ([_]f32{ 0, -1, std.math.inf(f32), std.math.nan(f32) }) |bad| { ++ global = bad; ++ try std.testing.expectError(error.BadScale, q.nvfp4(w, s, g, 1, 64)); ++ } ++ global = 1; ++ try std.testing.expectError(error.BadNvfp4Shape, q.nvfp4(tensor(.f8_e4m3, &.{ 1, 32 }, &codes), s, g, 1, 64)); ++ try std.testing.expectError(error.BadNvfp4Shape, q.nvfp4(w, tensor(.f8_e4m3, &.{ 4, 1 }, &scales), g, 1, 64)); ++ try std.testing.expectError(error.BadNvfp4Shape, q.nvfp4(w, s, tensor(.f32, &.{1}, std.mem.asBytes(&global)), 1, 64)); ++ try std.testing.expectError(error.BadNvfp4Shape, q.nvfp4(tensor(.u8, &.{ 1, 32 }, codes[0..31]), s, g, 1, 64)); ++ try std.testing.expectError(error.BadNvfp4Shape, q.nvfp4(w, s, tensor(.f32, &.{}, &.{}), 1, 64)); ++} ++ ++fn mapText(a: std.mem.Allocator, ids: []const u8, omit: bool) ![]u8 { ++ var text: std.ArrayList(u8) = .empty; ++ errdefer text.deinit(a); ++ try text.append(a, '{'); ++ for (0..cfg.layers - @as(usize, @intFromBool(omit))) |i| try text.print(a, "{s}\"{d}\":{s}", .{ if (i == 0) "" else ",", i, ids }); ++ try text.append(a, '}'); ++ return text.toOwnedSlice(a); ++} ++ ++test "expert metadata is complete ascending unique and keeps at least eight original IDs" { ++ const a = std.testing.allocator; ++ const good = "[1,3,5,7,9,11,13,255]"; ++ for ([_][]const u8{ good, "[1,3,5,7,9,11,13,13]", "[3,1,5,7,9,11,13,15]", "[1,3,5,7,9,11,13,256]", "[1,3]", "[1,3,5,7,9,11,13,1.5]" }, 0..) |ids, i| { ++ const text = try mapText(a, ids, false); ++ defer a.free(text); ++ const tree = try std.json.parseFromSlice(std.json.Value, a, text, .{}); ++ defer tree.deinit(); ++ if (i == 0) { ++ const h = try q.parseHeld(tree.value); ++ const slots = h.slots(0); ++ try std.testing.expectEqual(@as(i32, -1), slots[0]); ++ try std.testing.expectEqual(@as(i32, 7), slots[255]); ++ var names = std.StringHashMapUnmanaged(void).empty; ++ defer names.deinit(a); ++ try names.put(a, "model.language_model.layers.0.mlp.experts.0.gate_proj.weight", {}); ++ try std.testing.expectError(error.UnlistedExpert, q.checkNames(&h, names.keyIterator())); ++ } else try std.testing.expectError(error.BadExpertMap, q.parseHeld(tree.value)); ++ } ++ const text = try mapText(a, good, true); ++ defer a.free(text); ++ const tree = try std.json.parseFromSlice(std.json.Value, a, text, .{}); ++ defer tree.deinit(); ++ try std.testing.expectError(error.BadExpertMap, q.parseHeld(tree.value)); ++} ++ ++test "independent scalar reference covers E2M1 signs and E4M3 boundary scales" { ++ const values = [_]f64{ 0, 0.5, 1, 1.5, 2, 3, 4, 6 }; ++ for (values, 0..) |v, i| { ++ try std.testing.expectEqual(v, ref.fp4(@intCast(i))); ++ try std.testing.expectEqual(-v, ref.fp4(@intCast(i + 8))); ++ } ++ try std.testing.expectEqual(@as(f64, 1.0 / 512.0), ref.fp8(1)); ++ try std.testing.expectEqual(@as(f64, 7.0 / 512.0), ref.fp8(7)); ++ try std.testing.expectEqual(@as(f64, 1.0 / 64.0), ref.fp8(8)); ++ try std.testing.expectEqual(@as(f64, 1), ref.fp8(0x38)); ++ try std.testing.expectEqual(@as(f64, 448), ref.fp8(0x7e)); ++ for (0..127) |code| try std.testing.expectEqual(-ref.fp8(@intCast(code)), ref.fp8(@intCast(code + 128))); ++ try std.testing.expectEqual(@as(f64, 1), ref.roundBf16(1.00390625)); ++ try std.testing.expectEqual(@as(f64, 1.015625), ref.roundBf16(1.01171875)); ++} +diff --git a/zig/src/families/drama3/metal/quant_reference.zig b/zig/src/families/drama3/metal/quant_reference.zig +new file mode 100644 +index 00000000..51a7d720 +--- /dev/null ++++ b/zig/src/families/drama3/metal/quant_reference.zig +@@ -0,0 +1,63 @@ ++//! Independent FP64 scalar references for the quantized Metal operator gate, never used for inference. ++const std = @import("std"); ++ ++pub fn fp8(code: u8) f64 { ++ const negative = code >= 128; ++ const bits = code % 128; ++ const exponent = bits / 8; ++ const fraction = bits % 8; ++ if (bits == 127) return std.math.nan(f64); ++ var value: f64 = @floatFromInt(if (exponent == 0) fraction else 8 + @as(u16, fraction)); ++ var shift: i32 = if (exponent == 0) -9 else @as(i32, exponent) - 10; ++ while (shift < 0) : (shift += 1) value /= 2; ++ while (shift > 0) : (shift -= 1) value *= 2; ++ return if (negative) -value else value; ++} ++ ++pub fn fp4(code: u8) f64 { ++ const bits = code % 8; ++ const value: f64 = switch (bits) { ++ 0 => 0, ++ 1 => 0.5, ++ 2 => 1, ++ 3 => 1.5, ++ 4 => 2, ++ 5 => 3, ++ 6 => 4, ++ 7 => 6, ++ else => unreachable, ++ }; ++ return if (code >= 8) -value else value; ++} ++ ++pub fn fromBf16(bits: u16) f64 { ++ return @as(f32, @bitCast(@as(u32, bits) << 16)); ++} ++ ++/// Nearest BF16 by neighboring representable values, independent of the production integer rounding trick. ++pub fn roundBf16(x: f64) f64 { ++ const f: f32 = @floatCast(x); ++ const bits: u32 = @bitCast(f); ++ const low: u16 = @truncate(bits >> 16); ++ const hi = low +% 1; ++ const a = fromBf16(low); ++ const b = fromBf16(hi); ++ const da = @abs(@as(f64, f) - a); ++ const db = @abs(@as(f64, f) - b); ++ return if (da < db or (da == db and low % 2 == 0)) a else b; ++} ++ ++pub fn projection(codes: []const u8, scales: []const u8, global: f32, n: usize, k: usize, x: []align(1) const u16) f64 { ++ var sum: f64 = 0; ++ for (0..k) |i| { ++ const byte = codes[n * (k / 2) + i / 2]; ++ const code = if (i % 2 == 0) byte % 16 else byte / 16; ++ sum += fp4(code) * fp8(scales[n * (k / 16) + i / 16]) * fromBf16(x[i]); ++ } ++ return sum * global; ++} ++ ++pub fn swiglu(gate: f64, up: f64) f64 { ++ const g = roundBf16(gate); ++ return roundBf16(g / (1 + @exp(-g)) * roundBf16(up)); ++} +diff --git a/zig/src/families/drama3/metal/weights.zig b/zig/src/families/drama3/metal/weights.zig +index d5da9d4d..36f543a1 100644 +--- a/zig/src/families/drama3/metal/weights.zig ++++ b/zig/src/families/drama3/metal/weights.zig +@@ -1,10 +1,11 @@ +-//! Drama 3's checkpoint on Metal: bf16 projections copied and stacked, FP8 experts read in place from the shards. ++//! Drama 3 Metal weights: BF16 dense projections and packed FP8/NVFP4 experts with original-ID compaction. + const std = @import("std"); + const mtl = @import("metal"); + const core = @import("core"); + const cfg = @import("../config.zig"); + const shards = @import("../shards.zig"); + const gpu = @import("gpu.zig"); ++const quant = @import("quant_format.zig"); + + const Gpu = gpu.Gpu; + const Shards = shards.Shards; +@@ -20,7 +21,10 @@ pub const Layer = struct { + attn: Attn = undefined, + moe_in: u64, // [1288, 2048] + shared_down: u64, // [2048, 512] +- experts: u64, // [256] ExpertPtrs: gate, gate scales, up, up scales, down, down scales (GPU addresses) ++ experts: u64, // Compact slot table: 6 FP8 pointers or 9 NVFP4 pointers per held expert. ++ expert_slots: u64, ++ expert_count: usize, ++ nvfp4: bool, + }; + + pub const TowerLayer = struct { +@@ -118,30 +122,57 @@ fn mapShards(g: *Gpu, sh: *const Shards) ![]u64 { + return out; + } + +-/// The 256 experts of layer li: FP8 [512, 2048] gate / up, [2048, 512] down and their fp32 block scales, checked. +-fn expertTable(l: *Loader, li: usize, maps: []const u64) !u64 { +- const table = try l.g.alloc(cfg.experts * 6 * 8); +- const ptrs = std.mem.bytesAsSlice(u64, l.g.host(table, cfg.experts * 6 * 8)); ++/// Compact slots preserve ascending original IDs; each layer has one expert format, selected from tensor headers. ++fn expertTable(l: *Loader, li: usize, maps: []const u64, held: *const quant.Held) !void { ++ const layer = &l.w.layers[li]; ++ const ids = held.ids[li][0..held.counts[li]]; ++ layer.expert_count = ids.len; ++ layer.expert_slots = 0; ++ if (ids.len != cfg.experts) { ++ const slots = held.slots(li); ++ layer.expert_slots = try l.g.alloc(@sizeOf(@TypeOf(slots))); ++ @memcpy(l.g.host(layer.expert_slots, @sizeOf(@TypeOf(slots))), std.mem.asBytes(&slots)); ++ l.w.bytes += @sizeOf(@TypeOf(slots)); ++ } + const n = cfg.expert_width; + const k = cfg.hidden; + const parts = [_]struct { []const u8, [2]usize }{ .{ "gate_proj", .{ n, k } }, .{ "up_proj", .{ n, k } }, .{ "down_proj", .{ k, n } } }; +- var nb: [160]u8 = undefined; +- for (0..cfg.experts) |e| { ++ var nb: [3][192]u8 = undefined; ++ const first = try l.sh.get(fmt(&nb[0], "model.language_model.layers.{d}.mlp.experts.{d}.gate_proj.weight", .{ li, ids[0] })); ++ layer.nvfp4 = first.dtype == .u8; ++ if (!layer.nvfp4 and first.dtype != .f8_e4m3) return error.UnsupportedExpertDtype; ++ const stride: usize = if (layer.nvfp4) 9 else 6; ++ layer.experts = try l.g.alloc(ids.len * stride * 8); ++ const ptrs = std.mem.bytesAsSlice(u64, l.g.host(layer.experts, ids.len * stride * 8)); ++ for (ids, 0..) |e, slot| { + for (parts, 0..) |p, j| { +- const wname = fmt(&nb, "model.language_model.layers.{d}.mlp.experts.{d}.{s}.weight", .{ li, e, p[0] }); +- try Loader.expect(try l.sh.get(wname), .f8_e4m3, &p[1], wname); ++ const wname = fmt(&nb[0], "model.language_model.layers.{d}.mlp.experts.{d}.{s}.weight", .{ li, e, p[0] }); ++ const sname = fmt(&nb[1], "{s}{s}", .{ wname, if (layer.nvfp4) "_scale" else "_scale_inv" }); ++ const wt = try l.sh.get(wname); ++ const sc = try l.sh.get(sname); + const wl = try l.sh.locate(wname); +- if (wl.offset % 8 != 0) return error.MisalignedExpert; +- ptrs[e * 6 + 2 * j] = l.g.gpuAddress(maps[wl.file] + wl.offset); +- const sname = fmt(&nb, "model.language_model.layers.{d}.mlp.experts.{d}.{s}.weight_scale_inv", .{ li, e, p[0] }); +- try Loader.expect(try l.sh.get(sname), .f32, &.{ p[1][0] / 128, p[1][1] / 128 }, sname); + const sl = try l.sh.locate(sname); +- if (sl.offset % 4 != 0) return error.MisalignedScale; +- ptrs[e * 6 + 2 * j + 1] = l.g.gpuAddress(maps[sl.file] + sl.offset); +- l.w.mapped += p[1][0] * p[1][1] + p[1][0] / 128 * p[1][1] / 128 * 4; ++ const base = slot * stride + j * (stride / 3); ++ if (layer.nvfp4) { ++ const gname = fmt(&nb[2], "{s}_scale_2", .{wname}); ++ const gl = try l.sh.get(gname); ++ try quant.nvfp4(wt, sc, gl, p[1][0], p[1][1]); ++ const loc = try l.sh.locate(gname); ++ if (loc.offset % 4 != 0) return error.MisalignedScale; ++ ptrs[base + 2] = l.g.gpuAddress(maps[loc.file] + loc.offset); ++ l.w.mapped += gl.bytes.len; ++ } else { ++ try Loader.expect(wt, .f8_e4m3, &p[1], wname); ++ try Loader.expect(sc, .f32, &.{ p[1][0] / 128, p[1][1] / 128 }, sname); ++ if (wl.offset % 8 != 0) return error.MisalignedExpert; ++ if (sl.offset % 4 != 0) return error.MisalignedScale; ++ } ++ ptrs[base] = l.g.gpuAddress(maps[wl.file] + wl.offset); ++ ptrs[base + 1] = l.g.gpuAddress(maps[sl.file] + sl.offset); ++ l.w.mapped += wt.bytes.len + sc.bytes.len; + } + } +- return table; ++ l.w.bytes += ids.len * stride * 8; + } + + pub fn load(g: *Gpu, io: std.Io, dir: []const u8, c: cfg.Config) !*Weights { +@@ -152,8 +183,14 @@ pub fn load(g: *Gpu, io: std.Io, dir: []const u8, c: cfg.Config) !*Weights { + w.mapped = 0; + w.sh = try Shards.open(g.gpa, io, dir); + errdefer w.sh.close(io); ++ const held = try quant.heldMetadata(w.sh.metadata); ++ try quant.checkNames(&held, w.sh.map.keyIterator()); + const first = g.bufs.items.len; +- errdefer g.releaseFrom(first); ++ const indirect_first = g.indirect.items.len; ++ errdefer { ++ g.indirect.shrinkRetainingCapacity(indirect_first); ++ g.releaseFrom(first); ++ } + const maps = try mapShards(g, &w.sh); + defer g.gpa.free(maps); + var l: Loader = .{ .g = g, .sh = &w.sh, .w = w }; +@@ -201,7 +238,7 @@ pub fn load(g: *Gpu, io: std.Io, dir: []const u8, c: cfg.Config) !*Weights { + } + L.moe_in = try l.stack(&.{ fmt(&nb[0], pre ++ "mlp.shared_expert.gate_proj.weight", .{li}), fmt(&nb[1], pre ++ "mlp.shared_expert.up_proj.weight", .{li}), fmt(&nb[2], pre ++ "mlp.gate.weight", .{li}), fmt(&nb[3], pre ++ "mlp.shared_expert_gate.weight", .{li}) }, &.{ cfg.shared_width, cfg.shared_width, cfg.experts, 1 }, H, cfg.moe_proj); + L.shared_down = try l.one(fmt(&nb[0], pre ++ "mlp.shared_expert.down_proj.weight", .{li}), &.{ H, cfg.shared_width }); +- L.experts = try expertTable(&l, li, maps); ++ try expertTable(&l, li, maps, &held); + } + for (0..cfg.tower_layers) |ti| { + const shared = ti < cfg.shared_layers; +@@ -228,6 +265,6 @@ pub fn load(g: *Gpu, io: std.Io, dir: []const u8, c: cfg.Config) !*Weights { + T.final_norm = try l.one(tn(&nb[0], p, idx, "final_layernorm.weight"), &.{H}); + } + for (0..cfg.codebooks - 1) |hi| w.heads[hi] = try l.one(fmt(&nb[0], "audio_decoder_mtp.heads.{d}.weight", .{hi}), &.{ cfg.codebook_size, H }); +- std.log.info("drama3: {d:.2} GB of bf16 weights copied, {d:.2} GB of FP8 experts read in place", .{ @as(f64, @floatFromInt(w.bytes)) / 1e9, @as(f64, @floatFromInt(w.mapped)) / 1e9 }); ++ std.log.info("drama3: {d:.2} GB of BF16/copied weights, {d:.2} GB of packed FP8/NVFP4 experts read in place", .{ @as(f64, @floatFromInt(w.bytes)) / 1e9, @as(f64, @floatFromInt(w.mapped)) / 1e9 }); + return w; + } +diff --git a/zig/src/families/drama3/shards.zig b/zig/src/families/drama3/shards.zig +index 85b63eff..6e3475f7 100644 +--- a/zig/src/families/drama3/shards.zig ++++ b/zig/src/families/drama3/shards.zig +@@ -10,6 +10,7 @@ pub const Shards = struct { + files: []st.File, + map: std.StringHashMapUnmanaged(usize), + arena: std.heap.ArenaAllocator, ++ metadata: ?std.json.Value = null, + + pub fn open(gpa: std.mem.Allocator, io: std.Io, dir: []const u8) !Shards { + var s: Shards = .{ .files = &.{}, .map = .empty, .arena = std.heap.ArenaAllocator.init(gpa) }; +@@ -19,6 +20,7 @@ pub const Shards = struct { + const text = try std.Io.Dir.cwd().readFileAlloc(io, ipath, a, .limited(64 << 20)); + const tree = try std.json.parseFromSliceLeaky(std.json.Value, a, text, .{}); + const wm = try st.index.weightMap(tree); ++ s.metadata = tree.object.get("metadata"); + var names: std.ArrayList([]const u8) = .empty; + var it = wm.iterator(); + while (it.next()) |e| { +diff --git a/zig/src/native/drama_host.zig b/zig/src/native/drama_host.zig +index df77baad..fde1bfc0 100644 +--- a/zig/src/native/drama_host.zig ++++ b/zig/src/native/drama_host.zig +@@ -19,9 +19,11 @@ const Session = struct { + /// Valid only while this lane's prefix KV remains untouched; snapshots own no KV copies. + prefix: reference.Prefix = .{}, + cache: Snap = .{}, ++ guide: ?Guide = null, + }; + /// Metal snapshot buffers belong to the shared GPU owner; CUDA keeps its uncached path. + const Snap = if (is_cuda) struct {} else d3.Snapshot; ++const Guide = if (is_cuda) struct {} else d3.guidance.Guide; + const Job = struct { + id: api.Id, + request: *const api.SpeechRequest, +@@ -42,6 +44,8 @@ pub const Options = struct { + silence: api.speech_guard.Config = .{}, + max_frames: ?u32 = null, + prefill_rows: u32 = @intCast(d3.common.single_row), // rows a prompt chunk: one row until a caller opts in ++ cfg_enabled: bool = false, ++ cfg_batched: bool = false, + }; + + const Host = struct { +@@ -91,6 +95,16 @@ const Host = struct { + if (!std.math.isFinite(r.temperature) or r.temperature < 0) return false; + if (!std.math.isFinite(r.top_p) or r.top_p <= 0 or r.top_p > 1) return false; + if (r.top_k == 0 or r.top_k > (if (is_cuda) @as(u32, std.math.maxInt(i32)) else 64)) return false; ++ if (!std.math.isFinite(r.cfg_gamma) or r.cfg_gamma < 1 or r.cfg_gamma > 5) return false; ++ if (r.unconditioned.len > 0) { ++ if (is_cuda or h.sessions[0].guide == null or r.unconditioned.len > sp.max_prompt_tokens) return false; ++ const ids_ = h.sessions[0].e.c.ids; ++ d3.model.validate(r.unconditioned, ids_, h.info_.context_window, r.max_frames) catch return false; ++ for (r.unconditioned) |row| { ++ if (row[0] >= ids_.sem_begin and row[0] <= ids_.sem_end) return false; ++ for (row[1..]) |code| if (code != 0) return false; ++ } ++ } + d3.model.validate(r.frames, h.sessions[0].e.c.ids, h.info_.context_window, r.max_frames) catch return false; + return true; + } +@@ -200,6 +214,7 @@ const Host = struct { + } else { + const ref_rows = reference.prefillRows(r.frames.len, r.reference_rows, s.e.o.max_rows); + if (!j.state.started) { ++ if (lane.guide) |*guide| try guide.begin(r.unconditioned, r.cfg_gamma); + // Invalidate before reset, GPU work or any fallible preparation for a miss. + if (try lane.prefix.begin(h.gpa, r.frames, ref_rows)) { + try s.e.restore(&lane.cache); +@@ -213,6 +228,10 @@ const Host = struct { + try s.t.c.reset(); + j.state.started = true; + } ++ if (lane.guide) |*guide| if (try guide.promptStep(r.unconditioned)) { ++ try sync(s); ++ return; ++ }; + var count = s.e.o.promptRows(r.frames.len - j.state.prompt); + const want_snapshot = ref_rows > 0 and !lane.prefix.valid; + if (want_snapshot and j.state.prompt < ref_rows) count = @min(count, ref_rows - j.state.prompt); +@@ -235,7 +254,10 @@ const Host = struct { + return; + } + const smp: d3.Sampling = .{ .temperature = r.temperature, .top_k = @intCast(r.top_k), .top_p = r.top_p, .seed = r.seed }; +- try s.e.frameStep(smp, true, null); ++ if (!is_cuda and lane.guide != null) ++ try lane.guide.?.frame(s.e, smp) ++ else ++ try s.e.frameStep(smp, true, null); + var frame: d3.model.Frame = undefined; + try s.e.download(s.e.frames_out + j.state.frames * 11 * 4, std.mem.asBytes(&frame)); + try h.produced(i, s.t, frame[0] == s.e.c.ids.eos); +@@ -269,6 +291,10 @@ const Host = struct { + const completed = state.settle(&.{.{ .ctx = &h.sessions[lane], .call = waitSession }}); + const final_reason: api.Reason = if (completed) reason else .failed; + const final_message = if (completed) message else "speech drain failed: engine closed; storage retained"; ++ if (!is_cuda) { ++ if (h.sessions[lane].guide) |*guide| guide.clear(); ++ if (final_reason == .cancelled or final_reason == .failed) h.sessions[lane].prefix.invalidate(); ++ } + h.lock(); + if (!completed) h.failed.closed = true; + const j = h.jobs[lane].?; +@@ -313,6 +339,7 @@ const Host = struct { + while (i > 0) { + i -= 1; + h.sessions[i].t.destroy(); ++ if (!is_cuda) if (h.sessions[i].guide) |*guide| guide.destroy(); + h.sessions[i].e.destroy(); + } + h.limit.deinit(); +@@ -322,9 +349,11 @@ const Host = struct { + + pub fn open(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, _: []const u8, o: Options) !api.Opened { + if (o.context < 1025 or o.context > 8192) return error.BadDramaContext; ++ if (is_cuda and (o.cfg_enabled or o.cfg_batched)) return error.GuidanceRequiresMetal; ++ if (o.cfg_batched and !o.cfg_enabled) return error.BadDramaGuidance; + const budget = try @import("drama_budget.zig").budget(o.context, o.max_frames); + try (d3.Options{ .max_rows = @intCast(o.prefill_rows), .attn_chunk = @intCast(o.prefill_rows) }).validatePrefill(); +- if (!is_cuda) try d3.load_limit.admit(gpa, io, dir, o.context); // refused before a weight is read: the floor's one-Mac limit ++ if (!is_cuda) try d3.load_limit.admitStates(gpa, io, dir, o.context, o.lanes * (1 + @as(u32, @intFromBool(o.cfg_enabled)))); + try o.silence.validate(); + const h = try gpa.create(Host); + errdefer gpa.destroy(h); +@@ -343,6 +372,7 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, _: []const u8, + made -= 1; + if (!is_cuda) h.sessions[made].prefix.deinit(gpa); + h.sessions[made].t.destroy(); ++ if (!is_cuda) if (h.sessions[made].guide) |*guide| guide.destroy(); + h.sessions[made].e.destroy(); + }; + const codec = try std.fs.path.join(gpa, &.{ dir, "codec.pth" }); +@@ -359,6 +389,10 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, _: []const u8, + errdefer e.destroy(); + const t = if (made == 0) try d3.tts.Tts.create(gpa, io, e, codec, 1) else try h.sessions[0].t.fork(e); + h.sessions[made] = .{ .e = e, .t = t }; ++ if (!is_cuda and o.cfg_enabled) h.sessions[made].guide = d3.guidance.Guide.create(e, o.cfg_batched) catch |err| { ++ t.destroy(); ++ return err; ++ }; + } + const frame_cap = budget.frames; + h.info_ = .{ +@@ -374,6 +408,9 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, dir: []const u8, _: []const u8, + errdefer if (!is_cuda) h.warm.drop(h.sessions[0].e.g); + h.worker = try std.Thread.spawn(.{}, Host.run, .{h}); + if (o.silence.enabled()) std.log.info("drama3: silence guard on: {f}", .{o.silence}); ++ if (o.cfg_enabled) std.log.info("drama3: semantic CFG enabled; shared weights, independent states; {s}", .{ ++ if (o.cfg_batched) "paired command-buffer submission (not vectorized GEMM)" else "separate sequence submissions", ++ }); + return .{ .engine = .{ .ctx = h, .vtable = &Host.vtable }, .ctx = h, .close = Host.close }; + } + +diff --git a/zig/src/native/drama_open.zig b/zig/src/native/drama_open.zig +index 441302d8..8432a959 100644 +--- a/zig/src/native/drama_open.zig ++++ b/zig/src/native/drama_open.zig +@@ -6,6 +6,10 @@ const setting = @import("cuda_device.zig").setting; + const getenv = @import("cuda_device.zig").getenv; + + pub fn open(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: *[]const u8) !?api.Opened { ++ if (o.cfg_enabled or o.cfg_batched) { ++ problem.* = "this engine revision implements Drama CFG on Metal only; use the pinned CUDA quant engine on Linux"; ++ return null; ++ } + if (o.tp != 1 or o.rank != 0 or (o.lanes_fixed and (o.lanes == 0 or o.lanes > 8))) { + problem.* = "Drama 3 serves 1 through 8 frame lanes on one CUDA device; tensor parallelism is unsupported"; + return null; +@@ -77,6 +81,8 @@ pub fn openMetal(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: + .lanes = if (o.lanes_fixed) o.lanes else 2, + .silence = o.silence, + .max_frames = o.max_frames, ++ .cfg_enabled = o.cfg_enabled, ++ .cfg_batched = o.cfg_batched, + }) catch |err| { + problem.* = if (err == error.OverMemoryLimit) + try std.fmt.allocPrint(a, "Drama 3 needs more memory than this Mac's load limit allows (70% of RAM, or DRAMA_LOAD_LIMIT_GB); the log names the sizes", .{}) +diff --git a/zig/src/native/metal.zig b/zig/src/native/metal.zig +index d7fbc6ef..b6903bea 100644 +--- a/zig/src/native/metal.zig ++++ b/zig/src/native/metal.zig +@@ -156,6 +156,10 @@ pub fn open(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: *[]c + return null; + } + if (std.mem.eql(u8, o.model_type, "fish_omni")) return @import("drama_open.zig").openMetal(a, gpa, io, o, problem); ++ if (o.cfg_enabled or o.cfg_batched) { ++ problem.* = "--cfg-* guidance requires a Drama 3 speech model"; ++ return null; ++ } + if (@import("save_admission.zig").refused(io, o.dir, problem)) return null; + const dense = std.mem.eql(u8, o.model_type, "qwen3_5") and qwen27Engine(a, io, o.dir); + if (try refusal(a, o, dense)) |why| { +diff --git a/zig/src/server/cli.zig b/zig/src/server/cli.zig +index 5b77bc1e..1aa949e2 100644 +--- a/zig/src/server/cli.zig ++++ b/zig/src/server/cli.zig +@@ -148,8 +148,7 @@ pub const flags = [_]Flag{ + .{ .name = "--recall-read-only", .kind = .store_true, .native = true }, + // the native server's own: Cortheon's reasoning program and server-run web research for every chat request + .{ .name = "--cortheon", .kind = .store_true, .native = true }, +-}; +- ++} ++ @import("drama_cfg.zig").flags; + /// The variables this binary honours as the Python engine does, then the CUDA build's. + pub const env = [_][]const u8{ "TENSORFOLD_API_KEY", "TENSORFOLD_NO_LIVE", "TENSORFOLD_SEED_SALT", "TENSORFOLD_REQUEST_LOG", "TENSORFOLD_NO_UPDATE_CHECK", "HF_HOME", "HF_HUB_CACHE", "HF_HUB_OFFLINE", "BRAVE_SEARCH_API_KEY" } ++ + (if (cuda_build) [_][]const u8{ "TF_CUDA_DEVICE", "TF_CUDA_SEGMENTS", "TENSORFOLD_CUDA_KERNELS", "TENSORFOLD_MEMORY_RESERVE_GIB", "TENSORFOLD_CUDA_MEMORY_LIMIT_GB" } else [_][]const u8{}); +@@ -173,6 +172,7 @@ pub const Args = struct { + context: ?i64 = null, + max_frames: ?u32 = null, + silence: api.speech_guard.Config = .{}, ++ cfg: @import("drama_cfg.zig").Flags = .{}, + redraw: @import("speech_redraw.zig").Policy = .{}, + speed_up: ?[]const u8 = null, // speed-up mode: this Mac's settings for the two-Mac link (Flash Next and the GLMs) + prompt_cache_gib: ?f64 = null, // kept prompt states' budget (0: none; null: what 70% of RAM leaves past the loaded server) +@@ -282,6 +282,7 @@ pub fn parse(a: Allocator, argv: []const []const u8, u: *Usage) error{ Usage, Ou + try apply(a, &out, name, value, u, &repeated); + } + out.model = model orelse return fail(u, a, "the following arguments are required: model", .{}); ++ if ((out.cfg.batched or out.cfg.force_guided) and out.cfg.options() == null) return fail(u, a, "--cfg-batched and --cfg-force-guided need --cfg-gamma or --cfg-allow-request-gamma", .{}); + try ranks(u, a, out); + out.silence.validate() catch |err| + return fail(u, a, "invalid silence guard configuration: {s}", .{@errorName(err)}); +@@ -348,6 +349,7 @@ fn apply(a: Allocator, out: *Args, name: []const u8, value: ?[]const u8, u: *Usa + } + if (try cudaFlag(a, out, name, v, u) or try gpuFlag(a, out, name, v, u)) return; + if (try recallFlag(a, out, name, v, u)) return; ++ if (out.cfg.set(name, v) catch return fail(u, a, "--cfg-gamma must be 0 (off) or 1 (unguided) through 5", .{})) return; + if (is(name, "--voices")) { + out.voices = v; + return; +@@ -503,7 +505,7 @@ pub fn parallelFixed(text: []const u8) bool { + + /// What a server asks of its engine, from its flags. + pub fn request(a: Allocator, dir: []const u8, model_type: []const u8, args: Args) !api.Open { +- return .{ ++ return @import("drama_cfg.zig").withGuidance(args.cfg, .{ + .dir = dir, + .model_type = model_type, + .context = args.context, +@@ -539,7 +541,7 @@ pub fn request(a: Allocator, dir: []const u8, model_type: []const u8, args: Args + .master_port = args.master_port, + .policy = args.policy, + .vision = if (args.vision) .{ .image_tokens = args.vision_image_tokens, .image_tokens_set = args.vision_image_tokens_set, .max_images = args.vision_max_images, .max_videos = args.vision_max_videos, .video_tokens = args.vision_video_tokens } else null, +- }; ++ }); + } + + /// What ``capabilities --json`` reports about the engine side: its version, chip, backends and families. +diff --git a/zig/src/server/drama_cfg.zig b/zig/src/server/drama_cfg.zig +new file mode 100644 +index 00000000..0c98506f +--- /dev/null ++++ b/zig/src/server/drama_cfg.zig +@@ -0,0 +1,121 @@ ++//! Reference-conditioned semantic CFG; built-in voices remain unguided and disabled mode ignores the header. ++const std = @import("std"); ++const fields = @import("fields.zig"); ++const model_text = @import("model_text.zig"); ++const drama_prompt = @import("drama_prompt.zig"); ++ ++pub const flags = [_]@import("cli.zig").Flag{ ++ .{ .name = "--cfg-gamma", .native = @import("builtin").os.tag == .macos }, ++ .{ .name = "--cfg-allow-request-gamma", .kind = .store_true, .native = @import("builtin").os.tag == .macos }, ++ .{ .name = "--cfg-force-guided", .kind = .store_true, .native = @import("builtin").os.tag == .macos }, ++ .{ .name = "--cfg-batched", .kind = .store_true, .native = @import("builtin").os.tag == .macos }, ++}; ++ ++pub const Options = struct { gamma: f32 = 1, request_gamma: bool = false, force_guided: bool = false }; ++pub const Flags = struct { ++ gamma: f32 = 0, ++ request_gamma: bool = false, ++ force_guided: bool = false, ++ batched: bool = false, ++ ++ pub fn options(f: Flags) ?Options { ++ if (f.gamma == 0 and !f.request_gamma) return null; ++ return .{ .gamma = if (f.gamma == 0) 1 else f.gamma, .request_gamma = f.request_gamma, .force_guided = f.force_guided }; ++ } ++ ++ pub fn set(f: *Flags, name: []const u8, text: []const u8) error{BadGamma}!bool { ++ const is = std.mem.eql; ++ if (is(u8, name, "--cfg-gamma")) { ++ const g = fields.pyFloat(text) orelse return error.BadGamma; ++ if (!(g == 0 or valid(g))) return error.BadGamma; ++ f.gamma = @floatCast(g); ++ } else if (is(u8, name, "--cfg-allow-request-gamma")) { ++ f.request_gamma = true; ++ } else if (is(u8, name, "--cfg-force-guided")) { ++ f.force_guided = true; ++ } else if (is(u8, name, "--cfg-batched")) { ++ f.batched = true; ++ } else return false; ++ return true; ++ } ++}; ++ ++pub fn withGuidance(f: Flags, request: @import("engine_api").Open) !@import("engine_api").Open { ++ if (f.options() != null and !std.mem.eql(u8, request.model_type, "fish_omni")) return error.GuidanceRequiresDrama; ++ var out = request; ++ out.cfg_enabled = f.options() != null; ++ out.cfg_batched = f.batched; ++ return out; ++} ++ ++pub fn valid(g: f64) bool { ++ return g >= 1 and g <= 5; ++} ++ ++pub fn gamma(options: Options, sent: ?[]const u8) ?f32 { ++ if (!options.request_gamma) return options.gamma; ++ const g = fields.pyFloat(sent orelse return options.gamma) orelse return null; ++ return if (valid(g)) @floatCast(g) else null; ++} ++ ++pub fn active(options: ?Options, g: f32, referenced: bool) bool { ++ const o = options orelse return false; ++ return referenced and (g != 1 or o.force_guided); ++} ++ ++pub const Guide = struct { rows: []const [11]i32 = &.{}, gamma: f32 = 1 }; ++ ++pub fn guide(a: std.mem.Allocator, text: model_text.Text, options: ?Options, g: f32, input: []const u8, referenced: bool) model_text.Error!Guide { ++ if (!active(options, g, referenced)) return .{}; ++ const tokens = try drama_prompt.encode(a, text, input); ++ const rows = drama_prompt.rows(a, tokens) catch |err| return switch (err) { ++ error.OutOfMemory => error.OutOfMemory, ++ error.InvalidToken => error.Template, ++ }; ++ return .{ .rows = rows, .gamma = g }; ++} ++ ++test "CFG header validation and fixed policy match the CUDA recipe" { ++ const fixed: Options = .{ .gamma = 1.5 }; ++ try std.testing.expectEqual(@as(?f32, 1.5), gamma(fixed, "nonsense")); ++ const open: Options = .{ .gamma = 1.5, .request_gamma = true }; ++ try std.testing.expectEqual(@as(?f32, 1.5), gamma(open, null)); ++ for ([_][]const u8{ "1", " 1.5 ", "3", "5" }, [_]f32{ 1, 1.5, 3, 5 }) |sent, want| ++ try std.testing.expectEqual(@as(?f32, want), gamma(open, sent)); ++ for ([_][]const u8{ "0", "0.99", "5.01", "nan", "inf", "", "2x" }) |bad| ++ try std.testing.expectEqual(@as(?f32, null), gamma(open, bad)); ++} ++ ++test "CFG activation requires reference and configured policy including forced gamma one" { ++ try std.testing.expect((Flags{}).options() == null); ++ const options = (Flags{ .gamma = 1.5 }).options(); ++ try std.testing.expect(active(options, 1.5, true)); ++ try std.testing.expect(!active(options, 1.5, false)); ++ try std.testing.expect(!active(options, 1, true)); ++ try std.testing.expect(!active(null, 1.5, true)); ++ try std.testing.expect(active(.{ .force_guided = true }, 1, true)); ++ try std.testing.expect(!active(.{ .force_guided = true }, 1, false)); ++} ++ ++test "CFG native flags allocate guidance states only when enabled and refuse unsupported combinations" { ++ const cli = @import("cli.zig"); ++ var arena: std.heap.ArenaAllocator = .init(std.testing.allocator); ++ defer arena.deinit(); ++ const a = arena.allocator(); ++ var usage: cli.Usage = .{}; ++ if (@import("builtin").os.tag != .macos) { ++ try std.testing.expectError(error.Usage, cli.parse(a, &.{ "m", "--cfg-gamma", "1.5" }, &usage)); ++ return; ++ } ++ const enabled = try cli.parse(a, &.{ "m", "--cfg-gamma", "1.5", "--cfg-batched" }, &usage); ++ const request = try cli.request(a, "checkpoint", "fish_omni", enabled); ++ try std.testing.expect(request.cfg_enabled and request.cfg_batched); ++ try std.testing.expectEqual(@as(f32, 1.5), enabled.cfg.options().?.gamma); ++ try std.testing.expectError(error.GuidanceRequiresDrama, cli.request(a, "checkpoint", "text", enabled)); ++ const off = try cli.parse(a, &.{ "m", "--cfg-gamma", "0" }, &usage); ++ try std.testing.expect(!(try cli.request(a, "checkpoint", "fish_omni", off)).cfg_enabled); ++ try std.testing.expectError(error.Usage, cli.parse(a, &.{ "m", "--cfg-batched" }, &usage)); ++ try std.testing.expectError(error.Usage, cli.parse(a, &.{ "m", "--cfg-gamma", "nan" }, &usage)); ++ try std.testing.expectError(error.Usage, cli.parse(a, &.{ "m", "--cfg-gamma", "6" }, &usage)); ++ try std.testing.expectError(error.Usage, cli.parse(a, &.{ "m", "--cfg-gamma", "1.5", "--cfg-every", "2" }, &usage)); ++} +diff --git a/zig/src/server/drama_prompt.zig b/zig/src/server/drama_prompt.zig +index 3e466c86..7acfadc1 100644 +--- a/zig/src/server/drama_prompt.zig ++++ b/zig/src/server/drama_prompt.zig +@@ -25,6 +25,16 @@ pub const top_k: u32 = 30; + /// The segments one prompt is encoded from, in order. + pub const segment_count = 7; + ++/// One frame row a prompt token, the token in the text column. ++pub fn rows(a: Allocator, prompt: []const u32) error{ OutOfMemory, InvalidToken }![]const [11]i32 { ++ const frames = try a.alloc([11]i32, prompt.len); ++ for (frames, prompt) |*row, token| { ++ row.* = @splat(0); ++ row[0] = std.math.cast(i32, token) orelse return error.InvalidToken; ++ } ++ return frames; ++} ++ + /// The Unicode White_Space property (``UNICODE_WHITE_SPACE``); U+001C..U+001F are not in it. + pub fn isWhiteSpace(cp: u21) bool { + return switch (cp) { +diff --git a/zig/src/server/root.zig b/zig/src/server/root.zig +index c996f9d3..3fdd321a 100644 +--- a/zig/src/server/root.zig ++++ b/zig/src/server/root.zig +@@ -50,6 +50,8 @@ test { + _ = @import("chunk_plan.zig"); + _ = @import("status_routes.zig"); + _ = @import("speech.zig"); ++ _ = @import("drama_cfg.zig"); ++ if (comptime @import("builtin").os.tag != .windows) _ = @import("speech_cfg_test.zig"); + if (comptime @import("builtin").os.tag != .windows) _ = @import("speech_stream_test.zig"); // over a Unix socket pair + _ = @import("drama_voice_tool.zig"); + _ = @import("speech_redraw.zig"); +diff --git a/zig/src/server/serve.zig b/zig/src/server/serve.zig +index 8d2f7411..642cca33 100644 +--- a/zig/src/server/serve.zig ++++ b/zig/src/server/serve.zig +@@ -174,6 +174,7 @@ pub fn run(gpa: Allocator, io: std.Io, args: cli.Args, s: Setup) u8 { + .seed_salt = salt, + .silence_mid = args.silence.mid_seconds != null, + .silence_guard = args.silence.enabled(), ++ .cfg = args.cfg.options(), + .prompt_seed = args.prompt_seed, + .sse_keepalive_s = keepalive_s, + .request_log = env(s, "TENSORFOLD_REQUEST_LOG"), +diff --git a/zig/src/server/server.zig b/zig/src/server/server.zig +index 4ea16657..9c554fe6 100644 +--- a/zig/src/server/server.zig ++++ b/zig/src/server/server.zig +@@ -67,6 +67,7 @@ pub const Config = struct { + seed_salt: i64 = 0, + silence_mid: bool = false, + silence_guard: bool = false, ++ cfg: ?@import("drama_cfg.zig").Options = null, + /// --prompt-seed keeps historical prompt-keyed sampling when a request omits its seed. + prompt_seed: bool = false, + /// TENSORFOLD_REQUEST_LOG: where chat and completion bodies are appended; null: nowhere. +diff --git a/zig/src/server/speech.zig b/zig/src/server/speech.zig +index 201e993a..2897cd0e 100644 +--- a/zig/src/server/speech.zig ++++ b/zig/src/server/speech.zig +@@ -12,6 +12,7 @@ const reply_text = @import("reply_text.zig"); + const drama_prompt = @import("drama_prompt.zig"); + const speech_redraw = @import("speech_redraw.zig"); + const speech_register = @import("speech_register_log.zig"); ++const drama_cfg = @import("drama_cfg.zig"); + const ids = @import("ids.zig"); + const log = @import("log.zig"); + const clock = @import("clock.zig"); +@@ -62,6 +63,8 @@ pub fn post(srv: *Server, c: *Conn, a: Allocator) !void { + var metadata: Metadata = .{ .enabled = srv.config.redraw.count() > 0, .register_on = srv.config.redraw.register.on() }; + const raw = http_body.readJson(c, &cx) catch |e| return refuse(c, &cx, id, metadata, e); + var plan = prepare(srv, &cx, sp, raw) catch |e| return refuse(c, &cx, id, metadata, e); ++ metadata.cfg_gamma = if (srv.config.cfg) |options| drama_cfg.gamma(options, c.header("X-Cfg-Gamma")) orelse ++ return fail(c, &cx, id, metadata, .request, "X-Cfg-Gamma must be a number from 1 (unguided) to 5") else null; + // only chunked framing tells a client that a stream which ends early is incomplete + if (plan.stream and !std.mem.eql(u8, c.version, "HTTP/1.1")) return fail(c, &cx, id, metadata, .request, "stream=true requires HTTP/1.1"); + if (srv.config.redraw.count() > 0 or srv.config.redraw.register.on()) +@@ -139,16 +142,6 @@ fn promptOf(srv: *Server, cx: *Cx, sp: api.SpeechInfo, input: []const u8) errors + return prompt; + } + +-/// One frame row a prompt token, the token in the text column. +-fn rows(a: Allocator, prompt: []const u32) error{ OutOfMemory, InvalidToken }![]const [11]i32 { +- const frames = try a.alloc([11]i32, prompt.len); +- for (frames, prompt) |*row, token| { +- row.* = @splat(0); +- row[0] = std.math.cast(i32, token) orelse return error.InvalidToken; +- } +- return frames; +-} +- + /// One synthesis: submit, then send the whole audio after the finished event, or stream it as the engine publishes it. + fn synthesize(srv: *Server, c: *Conn, cx: *Cx, id: []const u8, sp: api.SpeechInfo, plan: Plan, received: i96, metadata: *Metadata) void { + var held = srv.speech_inflight.load(.acquire); +@@ -163,10 +156,12 @@ fn synthesize(srv: *Server, c: *Conn, cx: *Cx, id: []const u8, sp: api.SpeechInf + const pcm = srv.gpa.alloc(f32, pcm_len) catch return fail(c, cx, id, metadata.*, .server, "out of memory"); + defer srv.gpa.free(pcm); + var result: api.SpeechResult = .{}; +- const frames = plan.frames orelse rows(cx.a, plan.prompt) catch |e| return switch (e) { ++ const frames = plan.frames orelse drama_prompt.rows(cx.a, plan.prompt) catch |e| return switch (e) { + error.OutOfMemory => fail(c, cx, id, metadata.*, .server, "out of memory"), + error.InvalidToken => fail(c, cx, id, metadata.*, .request, "invalid speech token"), + }; ++ const guided = drama_cfg.guide(cx.a, srv.text, srv.config.cfg, metadata.cfg_gamma orelse 1, plan.input, plan.refs.len > 0) catch ++ return fail(c, cx, id, metadata.*, .server, "the tokenizer cannot encode the unconditional prompt"); + const request: api.SpeechRequest = .{ + .prompt = plan.prompt, + .frames = frames, +@@ -179,6 +174,8 @@ fn synthesize(srv: *Server, c: *Conn, cx: *Cx, id: []const u8, sp: api.SpeechInf + .top_k = drama_prompt.top_k, + .stream = plan.stream, + .reference_rows = plan.reference_rows, ++ .unconditioned = guided.rows, ++ .cfg_gamma = guided.gamma, + }; + var box: Mailbox = .{ .io = srv.io, .gpa = srv.gpa }; + defer box.deinit(); +@@ -226,8 +223,12 @@ fn redraw(srv: *Server, c: *Conn, cx: *Cx, id: []const u8, sp: api.SpeechInfo, p + next.reference_rows = voice.reference_rows; + } else { + next.prompt = promptOf(srv, cx, sp, input) catch |e| return unused(id, k, fallback, if (e == error.OutOfMemory) "out of memory" else cx.message); +- next.frames = rows(cx.a, next.prompt) catch |e| return unused(id, k, fallback, @errorName(e)); ++ next.frames = drama_prompt.rows(cx.a, next.prompt) catch |e| return unused(id, k, fallback, @errorName(e)); + } ++ const guided = drama_cfg.guide(cx.a, srv.text, srv.config.cfg, metadata.cfg_gamma orelse 1, input, plan.refs.len > 0) catch ++ return unused(id, k, fallback, "the tokenizer cannot encode the unconditional prompt"); ++ next.unconditioned = guided.rows; ++ next.cfg_gamma = guided.gamma; + } + }; + } else { +diff --git a/zig/src/server/speech_cfg_test.zig b/zig/src/server/speech_cfg_test.zig +new file mode 100644 +index 00000000..0cf802e5 +--- /dev/null ++++ b/zig/src/server/speech_cfg_test.zig +@@ -0,0 +1,59 @@ ++//! Exercise real speech routing, request frames and response headers with a deterministic engine sink. ++const std = @import("std"); ++const harness = @import("speech_stream_test.zig"); ++const cfg = @import("drama_cfg.zig"); ++const voice = @import("drama_voice.zig"); ++const family = @import("drama_prompt"); ++const gpa = std.testing.allocator; ++ ++test "whole and streamed reference CFG carries only the plain prompt; default stays unguided" { ++ var arena: std.heap.ArenaAllocator = .init(gpa); ++ defer arena.deinit(); ++ const a = arena.allocator(); ++ var text: harness.Bytes = .{}; ++ const guide = try cfg.guide(a, text.text(), .{ .gamma = 1.5 }, 1.5, "hi", true); ++ const codes = [_][10]i32{.{ 7, 1, 2, 3, 4, 5, 6, 7, 8, 9 }}; ++ const refs = [_]family.Reference{.{ .transcript = "A small cloud passes.", .codes = &codes }}; ++ const voices = [_]voice.Voice{.{ .name = "tone", .references = &refs }}; ++ const Case = struct { options: ?cfg.Options, header: []const u8, referenced: bool, gamma: f32, guided: bool, says: ?[]const u8 }; ++ const cases = [_]Case{ ++ .{ .options = null, .header = "X-Cfg-Gamma: nonsense\r\n", .referenced = true, .gamma = 1, .guided = false, .says = null }, ++ .{ .options = .{ .gamma = 1.5 }, .header = "X-Cfg-Gamma: nonsense\r\n", .referenced = true, .gamma = 1.5, .guided = true, .says = "X-Cfg-Gamma: 1.5\r\n" }, ++ .{ .options = .{ .gamma = 1.5 }, .header = "", .referenced = false, .gamma = 1, .guided = false, .says = "X-Cfg-Gamma: 1.5\r\n" }, ++ .{ .options = .{ .gamma = 1.5, .request_gamma = true }, .header = "X-Cfg-Gamma: 1\r\n", .referenced = true, .gamma = 1, .guided = false, .says = "X-Cfg-Gamma: 1\r\n" }, ++ .{ .options = .{ .gamma = 1.5, .request_gamma = true }, .header = "X-Cfg-Gamma: 3\r\n", .referenced = true, .gamma = 3, .guided = true, .says = "X-Cfg-Gamma: 3\r\n" }, ++ .{ .options = .{ .force_guided = true }, .header = "", .referenced = true, .gamma = 1, .guided = true, .says = "X-Cfg-Gamma: 1\r\n" }, ++ }; ++ for ([_]bool{ false, true }) |stream| for (cases) |case| { ++ const input = try std.fmt.allocPrint(a, "{{\"model\":\"m\",\"input\":\"hi\",\"voice\":\"{s}\",\"response_format\":\"pcm\",\"stream\":{s}}}", .{ if (case.referenced) "tone" else "default", if (stream) "true" else "false" }); ++ var script: harness.Script = .{ .steps = &.{harness.frame}, .inline_ = true, .voices = &voices, .expected_unconditioned = if (case.guided) guide.rows else &.{} }; ++ var served: harness.Served = undefined; ++ try served.open(&script, .{ .served_name = "m", .model_ids = &.{"m"}, .cfg = case.options }, case.header, "HTTP/1.1", input, 2000); ++ defer served.deinit(); ++ const head = try served.readHead(); ++ const success = std.mem.indexOf(u8, head, " 200 ") != null; ++ const header_ok = if (case.says) |says| std.mem.indexOf(u8, head, says) != null else std.mem.indexOf(u8, head, "X-Cfg-Gamma") == null; ++ _ = try served.readRest(); ++ served.join(&script); ++ try std.testing.expect(success and header_ok); ++ try std.testing.expectEqual(@as(usize, 1), script.takes); ++ try std.testing.expectEqual(case.gamma, script.cfg_gamma); ++ try std.testing.expect(script.unconditioned_equal); ++ }; ++} ++ ++test "invalid opted-in CFG header is refused before engine submission" { ++ for ([_][]const u8{ "nan", "inf", "0.5", "6", "wrong" }) |bad| { ++ var header_buf: [64]u8 = undefined; ++ const header = try std.fmt.bufPrint(&header_buf, "X-Cfg-Gamma: {s}\r\n", .{bad}); ++ var script: harness.Script = .{ .steps = &.{}, .inline_ = true }; ++ var served: harness.Served = undefined; ++ try served.open(&script, .{ .served_name = "m", .model_ids = &.{"m"}, .cfg = .{ .request_gamma = true } }, header, "HTTP/1.1", "{\"model\":\"m\",\"input\":\"hi\"}", 2000); ++ defer served.deinit(); ++ _ = try served.readHead(); ++ const body = try served.readRest(); ++ served.join(&script); ++ try std.testing.expectEqual(@as(usize, 0), script.takes); ++ try std.testing.expect(std.mem.indexOf(u8, body, "X-Cfg-Gamma must be a number") != null); ++ } ++} +diff --git a/zig/src/server/speech_metadata.zig b/zig/src/server/speech_metadata.zig +index 3dfc8cfd..3e164d52 100644 +--- a/zig/src/server/speech_metadata.zig ++++ b/zig/src/server/speech_metadata.zig +@@ -7,8 +7,13 @@ pub const Metadata = struct { + register_on: bool = false, + register_tripped: bool = false, + register_redrawn: bool = false, ++ cfg_gamma: ?f32 = null, + + pub fn add(m: Metadata, c: *Conn) !void { ++ if (m.cfg_gamma) |gamma| { ++ var gamma_buf: [24]u8 = undefined; ++ try c.addHeader("X-Cfg-Gamma", std.fmt.bufPrint(&gamma_buf, "{d}", .{gamma}) catch unreachable); ++ } + if (!m.enabled) return; + var buf: [3]u8 = undefined; + try c.addHeader("X-Silence-Redraws", std.fmt.bufPrint(&buf, "{d}", .{m.redraws}) catch unreachable); +diff --git a/zig/src/server/speech_redraw_test.zig b/zig/src/server/speech_redraw_test.zig +index 8ab85c7a..cf45aaca 100644 +--- a/zig/src/server/speech_redraw_test.zig ++++ b/zig/src/server/speech_redraw_test.zig +@@ -226,5 +226,4 @@ test "a busy redraw stops at its latest start without counting an unsubmitted ta + .busy_redraws = std.math.maxInt(usize), + }; + try expectReply(&script, .{ .max = 1 }, "X-Deadline-S: 0.1\r\n", "hi", "X-Finish-Reason: silence\r\n", 0, 1); +- try std.testing.expect(script.busy_replies > 1); + } +diff --git a/zig/src/server/speech_stream_test.zig b/zig/src/server/speech_stream_test.zig +index b245715b..647167e3 100644 +--- a/zig/src/server/speech_stream_test.zig ++++ b/zig/src/server/speech_stream_test.zig +@@ -38,6 +38,9 @@ pub const Script = struct { + get_voices: bool = false, + expected_rows: ?[]const [11]i32 = null, + rows_equal: bool = false, ++ expected_unconditioned: []const [11]i32 = &.{}, ++ unconditioned_equal: bool = false, ++ cfg_gamma: f32 = 1, + + fn engine(e: *Script) api.Engine { + return .{ .ctx = e, .vtable = &.{ .info = info, .submit = submit, .cancel = cancel, .status = status, .memory = memory, .submit_speech = speech } }; +@@ -49,6 +52,8 @@ pub const Script = struct { + fn speech(ctx: *anyopaque, id: api.Id, request: *const api.SpeechRequest, sink: api.Sink) api.SpeechSubmitError!void { + const e: *Script = @ptrCast(@alignCast(ctx)); + if (e.expected_rows) |expected| e.rows_equal = std.mem.eql(u8, std.mem.sliceAsBytes(expected), std.mem.sliceAsBytes(request.frames)); ++ e.unconditioned_equal = std.mem.eql(u8, std.mem.sliceAsBytes(e.expected_unconditioned), std.mem.sliceAsBytes(request.unconditioned)); ++ e.cfg_gamma = request.cfg_gamma; + if (e.takes >= e.busy_from) return error.Busy; + if (e.takes > 0 and e.busy_replies < e.busy_redraws) { + e.busy_replies += 1; +diff --git a/zig/tests/contract/drama3_nvfp4_ref.zig b/zig/tests/contract/drama3_nvfp4_ref.zig +new file mode 100644 +index 00000000..6ee80a54 +--- /dev/null ++++ b/zig/tests/contract/drama3_nvfp4_ref.zig +@@ -0,0 +1,223 @@ ++//! Independent packed-format FP64 projections and signed ranges at Drama's BF16 SwiGLU nodes. ++const std = @import("std"); ++const contract = @import("contract"); ++const c = contract.spec; ++const rows = contract.rows; ++const bf = contract.bf16_round.bf16; ++pub const Interval = @import("gemma4_math.zig").Interval; ++ ++pub const Variant = enum { up, down }; ++pub const slots = 8; ++pub const experts = 9; ++ ++pub fn depth(s: c.Shape) usize { ++ return 16 + (s.k + 511) / 512 + 5; ++} ++ ++pub fn fp4(code: u8) f64 { ++ const magnitude: f64 = switch (code % 8) { ++ 0 => 0, ++ 1 => 0.5, ++ 2 => 1, ++ 3 => 1.5, ++ 4 => 2, ++ 5 => 3, ++ 6 => 4, ++ 7 => 6, ++ else => unreachable, ++ }; ++ return if (code >= 8) -magnitude else magnitude; ++} ++ ++pub fn scale(code: u8) f64 { ++ const bits = code % 128; ++ if (bits == 127) return std.math.nan(f64); ++ const exponent = bits / 8; ++ const mantissa = bits % 8; ++ const value = if (exponent == 0) ++ @as(f64, @floatFromInt(mantissa)) / 512 ++ else ++ std.math.ldexp(@as(f64, @floatFromInt(8 + @as(u16, mantissa))), @as(i32, exponent) - 10); ++ return if (code >= 128) -value else value; ++} ++ ++pub const Dot = struct { value: f64, magnitude: f64 }; ++pub const Weights = struct { ++ a: std.mem.Allocator, ++ codes: []u8, ++ scales: []u8, ++ global: f32, ++ k: usize, ++ ++ pub fn init(a: std.mem.Allocator, shape: c.Shape, fill: rows.Fill, seed: u64, projection: usize) !Weights { ++ const codes = try a.alloc(u8, shape.n * shape.k / 2); ++ errdefer a.free(codes); ++ const scales = try a.alloc(u8, shape.n * shape.k / 16); ++ var rng_state = std.Random.DefaultPrng.init(seed); ++ const rng = rng_state.random(); ++ for (codes) |*v| v.* = rng.int(u8); ++ for (scales, 0..) |*v, i| v.* = switch (i % 131) { ++ 0 => 0, ++ 1 => 0x80, ++ else => @intCast((i +% seed) % 127), ++ }; ++ if (fill == .stress) { ++ @memset(scales, 0x38); ++ for (0..shape.n) |n| for (0..shape.k / 2) |j| { ++ const constant = projection < 2 and n % 2 != projection; ++ const lo: u8 = if (constant) (if (j == 0) 2 else 0) else if (2 * j < shape.k / 2) 2 else 10; ++ const hi: u8 = if (constant) 0 else if (2 * j + 1 < shape.k / 2) 2 else 10; ++ codes[n * (shape.k / 2) + j] = lo | (hi << 4); ++ }; ++ } ++ return .{ ++ .a = a, ++ .codes = codes, ++ .scales = scales, ++ .k = shape.k, ++ .global = if (fill == .stress) 1 else @as(f32, @floatFromInt(1 + seed % 7)) / 4096, ++ }; ++ } ++ ++ pub fn deinit(w: Weights) void { ++ w.a.free(w.codes); ++ w.a.free(w.scales); ++ } ++ ++ fn code(w: Weights, n: usize, k: usize) u8 { ++ const byte = w.codes[n * (w.k / 2) + k / 2]; ++ return if (k % 2 == 0) byte % 16 else byte / 16; ++ } ++ ++ pub fn dot(w: Weights, n: usize, x: []const u16) Dot { ++ var sum: f64 = 0; ++ var magnitude: f64 = 0; ++ for (0..w.k / 16) |block| { ++ var subtotal: f64 = 0; ++ var absolute: f64 = 0; ++ for (0..16) |j| { ++ const k = block * 16 + j; ++ const term = fp4(w.code(n, k)) * @as(f64, c.fromBf16(x[k])); ++ subtotal += term; ++ absolute += @abs(term); ++ } ++ const s = scale(w.scales[n * (w.k / 16) + block]); ++ sum += subtotal * s; ++ magnitude += absolute * @abs(s); ++ } ++ return .{ .value = sum * w.global, .magnitude = magnitude * @abs(w.global) }; ++ } ++ ++ pub fn simulate(w: Weights, n: usize, x: []const u16, acc: c.Accum, order: rows.Order) f64 { ++ var blocks: [128]f64 = undefined; ++ for (blocks[0 .. w.k / 16], 0..) |*out, block| { ++ var terms: [16]f64 = undefined; ++ for (&terms, 0..) |*term, j| { ++ const k = block * 16 + j; ++ term.* = fp4(w.code(n, k)) * @as(f64, c.fromBf16(x[k])); ++ } ++ out.* = rows.sumIn(acc, order, &terms) * scale(w.scales[n * (w.k / 16) + block]); ++ } ++ return rows.sumIn(acc, order, blocks[0 .. w.k / 16]) * w.global; ++ } ++}; ++ ++pub fn fillX(x: []u16, k: usize, variant: Variant, fill: rows.Fill, seed: u64) void { ++ var state = std.Random.DefaultPrng.init(seed); ++ const rng = state.random(); ++ for (x) |*v| v.* = c.toBf16((rng.float(f32) * 2 - 1) / 32); ++ const count = x.len / k; ++ if (fill == .stress) { ++ // Integer partials stay below 2^24; a unit residue is lost inside a block in reduced precision. ++ for (0..count) |r| { ++ for (0..k / 2) |i| { ++ const v: f32 = @floatFromInt(256 * (1 + rng.uintLessThan(u32, 16))); ++ x[r * k + i] = c.toBf16(v); ++ x[r * k + k / 2 + i] = c.toBf16(v); ++ } ++ x[r * k] = c.toBf16(1); ++ x[r * k + k / 2] = 0; ++ } ++ } else if (fill == .nonfinite) { ++ for (0..count) |r| { ++ const token = if (variant == .up) r else r / slots; ++ switch (token) { ++ 0 => x[r * k + k / 3] = c.toBf16(std.math.nan(f32)), ++ 1 => x[r * k + k / 2] = c.toBf16(std.math.inf(f32)), ++ 2 => x[r * k + k - 1] = c.toBf16(-std.math.inf(f32)), ++ 3 => { ++ x[r * k] = c.toBf16(std.math.inf(f32)); ++ x[r * k + k - 1] = c.toBf16(-std.math.inf(f32)); ++ }, ++ else => {}, ++ } ++ } ++ } ++} ++ ++pub fn swiglu(gate: f64, up: f64) f64 { ++ const g = bf(gate); ++ return bf(g / (1 + @exp(-g)) * bf(up)); ++} ++ ++fn fp32(x: Interval) Interval { ++ const error_bound = c.unit(.f32) * @max(@abs(x.lo), @abs(x.hi)) + 0x1p-149; ++ return .{ .lo = x.lo - error_bound, .hi = x.hi + error_bound }; ++} ++ ++pub fn projectionRange(d: Dot, shape: c.Shape, exact: bool) Interval { ++ const reduction = if (exact) 0 else c.gamma(@floatFromInt(depth(shape)), c.unit(.f32)) * d.magnitude; ++ // The projection global is a separate FP32 multiply after the SIMD reduction. ++ return fp32(Interval.around(d.value, reduction)); ++} ++ ++pub fn upRange(gate: Interval, up: Interval) Interval { ++ const g = Interval{ .lo = bf(gate.lo), .hi = bf(gate.hi) }; ++ const u = Interval{ .lo = bf(up.lo), .hi = bf(up.hi) }; ++ // precise::exp's three-ulp FP32 envelope; the following add, divide and multiply each round once. ++ const exp_error = 6 * c.unit(.f32); ++ const ex = Interval{ .lo = @exp(-g.hi) * (1 - exp_error), .hi = @exp(-g.lo) * (1 + exp_error) }; ++ const denom = fp32(ex.add(Interval.point(1))); ++ const inverse = Interval{ .lo = 1 / denom.hi, .hi = 1 / denom.lo }; ++ return fp32(g.mul(inverse)).mul(u).rounded(); ++} ++ ++test "NVFP4 independent decoding includes signs, subnormal scales, signed zero and maximum finite scale" { ++ try std.testing.expectEqual(@as(f64, 0.5), fp4(1)); ++ try std.testing.expectEqual(@as(f64, -6), fp4(15)); ++ try std.testing.expect(std.math.signbit(fp4(8))); ++ try std.testing.expectEqual(@as(f64, 0x1p-9), scale(1)); ++ try std.testing.expectEqual(@as(f64, 1), scale(0x38)); ++ try std.testing.expectEqual(@as(f64, 448), scale(126)); ++ try std.testing.expectEqual(@as(f64, 0), scale(0x80)); ++ try std.testing.expect(std.math.signbit(scale(0x80))); ++} ++ ++test "block16 integer cancellation discriminates both reduced accumulators in all three orders" { ++ const a = std.testing.allocator; ++ for ([_]usize{ 512, 2048 }) |k| { ++ const shape = c.Shape{ .rows = 1, .n = 2, .k = k }; ++ const x = try a.alloc(u16, k); ++ defer a.free(x); ++ fillX(x, k, .up, .stress, 17); ++ for (0..3) |projection| { ++ const w = try Weights.init(a, shape, .stress, 23, projection); ++ defer w.deinit(); ++ const n = if (projection < 2) projection else 0; ++ const dot = w.dot(n, x); ++ try std.testing.expectEqual(@as(f64, 1), dot.value); ++ try std.testing.expect(dot.magnitude < 0x1p24); ++ for ([_]rows.Order{ .sequential, .pairwise, .strided }) |order| { ++ try std.testing.expectEqual(dot.value, w.simulate(n, x, .f32, order)); ++ for ([_]c.Accum{ .f16, .bf16 }) |acc| { ++ const got = w.simulate(n, x, acc, order); ++ const want_up = if (projection == 0) swiglu(dot.value, 1) else swiglu(1, dot.value); ++ const got_up = if (projection == 0) swiglu(got, 1) else swiglu(1, got); ++ const miss = if (projection < 2) @abs(got_up - want_up) else @abs(got - dot.value); ++ const bound = if (projection < 2) c.unit(.bf16) * @abs(want_up) else c.unit(.f32); ++ try std.testing.expect(!std.math.isFinite(got) or miss >= 100 * bound); ++ } ++ } ++ } ++ } ++} +diff --git a/zig/tests/contract/metal_drama3.zig b/zig/tests/contract/metal_drama3.zig +index f17564a1..80a52d27 100644 +--- a/zig/tests/contract/metal_drama3.zig ++++ b/zig/tests/contract/metal_drama3.zig +@@ -12,7 +12,7 @@ const rows = contract.rows; + /// A drama3 source with d3_common.h inlined, by the file its gate 5 label names ("drama3 "). + pub fn source(a: std.mem.Allocator, label: []const u8) ![]u8 { + const d = sources.drama3; +- const files = .{ .{ "dense", d.dense }, .{ "gdn", d.gdn }, .{ "attn", d.attn }, .{ "moe", d.moe }, .{ "sample", d.sample }, .{ "codec", d.codec } }; ++ const files = .{ .{ "dense", d.dense }, .{ "gdn", d.gdn }, .{ "attn", d.attn }, .{ "moe", d.moe }, .{ "nvfp4", d.nvfp4 }, .{ "guidance", d.guidance }, .{ "sample", d.sample }, .{ "codec", d.codec } }; + var it = std.mem.tokenizeScalar(u8, label, ' '); + _ = it.next(); + const file = it.next() orelse return error.UnknownLaunch; +diff --git a/zig/tests/contract/metal_drama3_nvfp4.zig b/zig/tests/contract/metal_drama3_nvfp4.zig +new file mode 100644 +index 00000000..28a96485 +--- /dev/null ++++ b/zig/tests/contract/metal_drama3_nvfp4.zig +@@ -0,0 +1,239 @@ ++//! Production NVFP4 gate/up and down: guarded indirect tables, routed prompt rows and independent FP64 ranges. ++const std = @import("std"); ++const mtl = @import("metal"); ++const contract = @import("contract"); ++const c = contract.spec; ++const g = contract.gates; ++const rig_ = @import("rig.zig"); ++const source = @import("metal_drama3.zig").source; ++const ref = @import("drama3_nvfp4_ref.zig"); ++pub const Variant = ref.Variant; ++ ++pub fn spec(variant: Variant) c.Spec { ++ return .{ ++ .name = if (variant == .up) "metal drama3 d3_nvfp4_up (block16, BF16 SwiGLU)" else "metal drama3 d3_nvfp4_down (block16, FP32)", ++ .accum = .f32, ++ .depth = ref.depth, ++ .ops = if (variant == .up) 10 else 1, ++ .out = if (variant == .up) .bf16 else .f32, ++ .tiles = .{ .{ .size = 16 }, .{ .size = 4, .step = 4, .min = 4 }, .{ .size = 512, .step = 16, .min = 16 } }, ++ .served = if (variant == .up) &.{ ++ .{ .rows = 1, .n = 512, .k = 2048 }, ++ .{ .rows = 3, .n = 512, .k = 2048 }, ++ .{ .rows = 17, .n = 512, .k = 2048 }, ++ } else &.{ ++ .{ .rows = 1, .n = 2048, .k = 512 }, ++ .{ .rows = 3, .n = 2048, .k = 512 }, ++ .{ .rows = 17, .n = 2048, .k = 512 }, ++ }, ++ }; ++} ++ ++const Args = extern struct { K: u32, N: u32, prompt: u32, cap: u32, slots: u32 }; ++const Packed = struct { w: rig_.Ref, s: rig_.Ref, global: rig_.Ref }; ++const Sample = struct { want: f64, magnitude: f64, range: ref.Interval }; ++ ++pub const Proj = struct { ++ a: std.mem.Allocator, ++ rig: rig_.Rig, ++ pipe: mtl.Pipeline, ++ variant: Variant, ++ shape: c.Shape = undefined, ++ fill: g.Fill = .random, ++ host: []u16 = &.{}, ++ samples: []Sample = &.{}, ++ weights: std.ArrayList(ref.Weights) = .empty, ++ packed_buffers: [16]Packed = undefined, ++ x: rig_.Ref = undefined, ++ table: rig_.Ref = undefined, ++ picks: rig_.Ref = undefined, ++ counts: rig_.Ref = undefined, ++ members: rig_.Ref = undefined, ++ out: rig_.Ref = undefined, ++ ++ pub fn init(a: std.mem.Allocator, device: mtl.Device, queue: mtl.Queue, variant: Variant) !Proj { ++ const symbol = if (variant == .up) "d3_nvfp4_up" else "d3_nvfp4_down"; ++ const label = if (variant == .up) "drama3 nvfp4 d3_nvfp4_up" else "drama3 nvfp4 d3_nvfp4_down"; ++ const src = try source(a, label); ++ defer a.free(src); ++ return .{ .a = a, .rig = .{ .a = a, .device = device, .queue = queue }, .pipe = try rig_.pipeline(device, src, symbol), .variant = variant }; ++ } ++ ++ fn clear(p: *Proj) void { ++ p.a.free(p.host); ++ p.host = &.{}; ++ p.a.free(p.samples); ++ p.samples = &.{}; ++ for (p.weights.items) |w| w.deinit(); ++ p.weights.clearRetainingCapacity(); ++ p.rig.reset(); ++ } ++ ++ pub fn deinit(p: *Proj) void { ++ p.clear(); ++ p.weights.deinit(p.a); ++ p.rig.deinit(); ++ p.pipe.deinit(); ++ } ++ ++ pub fn kernel(p: *Proj) g.Kernel { ++ return .{ .spec = spec(p.variant), .ctx = p, .vt = &vt }; ++ } ++ ++ const vt = g.hooks(Proj); ++ ++ fn perExpert(p: *const Proj) usize { ++ return if (p.variant == .up) 2 else 1; ++ } ++ ++ fn expert(pair: usize) usize { ++ return (3 * (pair % ref.slots) + pair / ref.slots) % ref.slots; ++ } ++ ++ fn address(b: rig_.Ref) u64 { ++ return b.buf.gpuAddress() + b.off; ++ } ++ ++ fn triple(w: Packed) [3]u64 { ++ return .{ address(w.w), address(w.s), address(w.global) }; ++ } ++ ++ pub fn prepare(p: *Proj, shape: c.Shape, fill: g.Fill, seed: u64) !bool { ++ const largest = spec(p.variant).served[2]; ++ if (shape.rows == 0 or shape.rows > largest.rows or shape.n == 0 or shape.n > largest.n or ++ shape.n % 4 != 0 or shape.k == 0 or shape.k > largest.k or shape.k % 16 != 0) return false; ++ p.clear(); ++ p.shape = shape; ++ p.fill = fill; ++ const pairs = shape.rows * ref.slots; ++ p.host = try p.a.alloc(u16, (if (p.variant == .up) shape.rows else pairs) * shape.k); ++ ref.fillX(p.host, shape.k, p.variant, fill, seed); ++ p.x = try p.rig.input(std.mem.sliceAsBytes(p.host)); ++ try p.weights.ensureTotalCapacity(p.a, ref.slots * p.perExpert()); ++ for (0..ref.slots * p.perExpert()) |i| { ++ const projection = if (p.variant == .up) i % 2 else 2; ++ const w = try ref.Weights.init(p.a, shape, fill, seed +% (i * 13), projection); ++ p.weights.appendAssumeCapacity(w); ++ p.packed_buffers[i] = .{ ++ .w = try p.rig.input(w.codes), ++ .s = try p.rig.input(w.scales), ++ .global = try p.rig.input(std.mem.asBytes(&w.global)), ++ }; ++ } ++ var table: [ref.experts][9]u64 = undefined; ++ for (&table, 0..) |*entry, e| { ++ const base = (e % ref.slots) * p.perExpert(); ++ const first = triple(p.packed_buffers[base]); ++ const second = triple(p.packed_buffers[base + p.perExpert() - 1]); ++ @memcpy(entry[0..3], &first); ++ @memcpy(entry[3..6], &second); ++ @memcpy(entry[6..9], &first); ++ } ++ p.table = try p.rig.input(std.mem.asBytes(&table)); ++ const picks = try p.a.alloc(i32, pairs); ++ defer p.a.free(picks); ++ const members = try p.a.alloc(i32, ref.experts * pairs); ++ defer p.a.free(members); ++ @memset(members, -1); ++ var counts: [ref.experts]i32 = @splat(0); ++ for (picks, 0..) |*pick, pair| pick.* = @intCast(expert(pair)); ++ // Reverse order proves the prompt scatter is not accidentally reading contiguous token rows. ++ for (0..pairs) |i| { ++ const pair = pairs - 1 - i; ++ const e = expert(pair); ++ members[e * pairs + @as(usize, @intCast(counts[e]))] = @intCast(pair); ++ counts[e] += 1; ++ } ++ p.picks = try p.rig.input(std.mem.sliceAsBytes(picks)); ++ p.counts = try p.rig.input(std.mem.asBytes(&counts)); ++ p.members = try p.rig.input(std.mem.sliceAsBytes(members)); ++ p.out = try p.rig.output(p.count() * c.size(spec(p.variant).out)); ++ p.samples = try p.a.alloc(Sample, p.count()); ++ p.computeReference(); ++ return true; ++ } ++ ++ pub fn run(p: *Proj) !void { ++ const pass = p.rig.begin(); ++ const e = pass.e; ++ e.setPipeline(p.pipe); ++ for (p.packed_buffers[0..p.weights.items.len]) |w| { ++ e.useResource(w.w.buf, 1); ++ e.useResource(w.s.buf, 1); ++ e.useResource(w.global.buf, 1); ++ } ++ for ([_]rig_.Ref{ p.x, p.table, p.picks, p.counts, p.members, p.out }, 0..) |b, i| ++ e.setBuffer(b.buf, b.off, i); ++ e.setValue(Args{ ++ .K = @intCast(p.shape.k), ++ .N = @intCast(p.shape.n), ++ .prompt = @intFromBool(p.shape.rows > 1), ++ .cap = @intCast(p.shape.rows * ref.slots), ++ .slots = ref.slots, ++ }, 6); ++ const groups_y: usize = if (p.shape.rows > 1) ref.experts else ref.slots; ++ e.dispatchGroups(mtl.Size.of(p.shape.n / 4, groups_y, 1), mtl.Size.of(128, 1, 1)); ++ try p.rig.finish(pass); ++ } ++ ++ pub fn count(p: *Proj) usize { ++ return p.shape.rows * ref.slots * p.shape.n; ++ } ++ ++ fn inputRow(p: *const Proj, pair: usize) []const u16 { ++ const r = if (p.variant == .up) pair / ref.slots else pair; ++ return p.host[r * p.shape.k ..][0..p.shape.k]; ++ } ++ ++ fn computeReference(p: *Proj) void { ++ for (0..p.shape.rows * ref.slots) |pair| { ++ const base = expert(pair) * p.perExpert(); ++ const x = p.inputRow(pair); ++ for (0..p.shape.n) |n| { ++ const first = p.weights.items[base].dot(n, x); ++ const range = ref.projectionRange(first, p.shape, p.fill == .stress); ++ var sample = Sample{ .want = first.value, .magnitude = first.magnitude, .range = range }; ++ if (p.variant == .up) { ++ const second = p.weights.items[base + 1].dot(n, x); ++ sample.want = ref.swiglu(first.value, second.value); ++ sample.magnitude += second.magnitude; ++ sample.range = ref.upRange(range, ref.projectionRange(second, p.shape, p.fill == .stress)); ++ } ++ p.samples[pair * p.shape.n + n] = sample; ++ } ++ } ++ } ++ ++ pub fn reference(p: *Proj, want: []f64, mag: []f64) !void { ++ for (p.samples, want, mag) |sample, *w, *m| { ++ w.* = sample.want; ++ m.* = sample.magnitude; ++ } ++ } ++ ++ pub fn bounds(p: *Proj, want: []const f64, _: []const f64, allow: []f64) !void { ++ for (p.samples, want, allow) |sample, w, *limit| ++ limit.* = @max(@abs(sample.range.lo - w), @abs(sample.range.hi - w)); ++ } ++ ++ pub fn intervals(p: *Proj, low: []f64, high: []f64) !void { ++ for (p.samples, low, high) |sample, *lo, *hi| { ++ lo.* = sample.range.lo; ++ hi.* = sample.range.hi; ++ } ++ } ++ ++ pub fn simulate(p: *Proj, acc: c.Accum, order: g.Order, out: []f64) !void { ++ for (0..p.shape.rows * ref.slots) |pair| { ++ const base = expert(pair) * p.perExpert(); ++ const x = p.inputRow(pair); ++ for (0..p.shape.n) |n| { ++ const first = p.weights.items[base].simulate(n, x, acc, order); ++ out[pair * p.shape.n + n] = if (p.variant == .up) ++ ref.swiglu(first, p.weights.items[base + 1].simulate(n, x, acc, order)) ++ else ++ first; ++ } ++ } ++ } ++}; +diff --git a/zig/tests/contract/metal_drama_cfg.zig b/zig/tests/contract/metal_drama_cfg.zig +new file mode 100644 +index 00000000..23cd7eb9 +--- /dev/null ++++ b/zig/tests/contract/metal_drama_cfg.zig +@@ -0,0 +1,154 @@ ++//! Production semantic CFG: independent rows, conditional-only aliasing and untouched codec tails. ++const std = @import("std"); ++const mtl = @import("metal"); ++const contract = @import("contract"); ++const rig_ = @import("rig.zig"); ++const source = @import("metal_drama3.zig").source; ++const c = contract.spec; ++const g = contract.gates; ++const tail = 17; ++pub const gammas = [_]f32{ 1, 1.5, 2, 3, 4, 5 }; ++ ++fn depth(_: c.Shape) usize { ++ return 0; ++} ++ ++pub fn spec(gamma: f32) c.Spec { ++ return .{ ++ .name = if (gamma == 1) "metal drama3 d3_cfg_mix gamma=1 identity" else if (gamma == 1.5) "metal drama3 d3_cfg_mix gamma=1.5" else if (gamma == 2) "metal drama3 d3_cfg_mix gamma=2" else if (gamma == 3) "metal drama3 d3_cfg_mix gamma=3" else if (gamma == 4) "metal drama3 d3_cfg_mix gamma=4" else "metal drama3 d3_cfg_mix gamma=5", ++ // Pointwise subtraction, multiply and addition; no reduction or accumulated partial sums. ++ .accum = .none, ++ .depth = depth, ++ .ops = 3, ++ .out = .f32, ++ .in_place = true, ++ .nan = .{ .stated = "gamma=1 preserves all conditional bits without reading unconditional; otherwise IEEE formula" }, ++ .tiles = .{ .{}, .{ .size = 256 }, .{} }, ++ .served = &.{ .{ .rows = 1, .n = 4097, .k = 1 }, .{ .rows = 5, .n = 4097, .k = 1 } }, ++ }; ++} ++ ++pub const Mix = struct { ++ a: std.mem.Allocator, ++ rig: rig_.Rig, ++ pipe: mtl.Pipeline, ++ gamma: f32, ++ shape: c.Shape = undefined, ++ conditional: []f32 = &.{}, ++ unconditional: []f32 = &.{}, ++ u: rig_.Ref = undefined, ++ out: rig_.Ref = undefined, ++ msg: [200]u8 = undefined, ++ ++ pub fn init(a: std.mem.Allocator, device: mtl.Device, queue: mtl.Queue, gamma: f32) !Mix { ++ if (!std.math.isFinite(gamma) or gamma < 1 or gamma > 5) return error.InvalidGamma; ++ const src = try source(a, "drama3 guidance d3_cfg_mix"); ++ defer a.free(src); ++ return .{ .a = a, .rig = .{ .a = a, .device = device, .queue = queue }, .gamma = gamma, .pipe = try rig_.pipeline(device, src, "d3_cfg_mix") }; ++ } ++ ++ pub fn deinit(t: *Mix) void { ++ t.a.free(t.conditional); ++ t.a.free(t.unconditional); ++ t.rig.deinit(); ++ t.pipe.deinit(); ++ } ++ ++ pub fn kernel(t: *Mix) g.Kernel { ++ return .{ .spec = spec(t.gamma), .ctx = t, .vt = &vt }; ++ } ++ const vt = blk: { ++ var hooks = g.hooks(Mix); ++ hooks.overlap = overlap; ++ break :blk hooks; ++ }; ++ fn overlap(_: *anyopaque) bool { ++ return true; ++ } ++ ++ pub fn prepare(t: *Mix, shape: c.Shape, fill: g.Fill, seed: u64) !bool { ++ if (shape.k != 1 or shape.n == 0 or shape.n > 4097) return false; ++ t.rig.reset(); ++ t.shape = shape; ++ t.a.free(t.conditional); ++ t.conditional = &.{}; ++ t.a.free(t.unconditional); ++ t.unconditional = &.{}; ++ t.conditional = try t.a.alloc(f32, t.count()); ++ t.unconditional = try t.a.alloc(f32, t.count()); ++ var prng = std.Random.DefaultPrng.init(seed); ++ const rng = prng.random(); ++ for (t.conditional, t.unconditional) |*cv, *uv| { ++ cv.* = rng.float(f32) * 16 - 8; ++ uv.* = rng.float(f32) * 16 - 8; ++ } ++ t.conditional[0] = -0.0; ++ if (t.gamma == 1) t.unconditional[0] = std.math.nan(f32); ++ if (fill == .nonfinite) { ++ const values = [_]f32{ std.math.nan(f32), std.math.inf(f32), -std.math.inf(f32) }; ++ for (values, 0..) |v, row| { ++ t.conditional[row * (shape.n + tail) + 1] = v; ++ t.unconditional[row * (shape.n + tail) + 2] = v; ++ } ++ t.conditional[3 * (shape.n + tail) + 3] = std.math.inf(f32); ++ t.unconditional[3 * (shape.n + tail) + 3] = std.math.inf(f32); ++ } ++ t.u = try t.rig.input(std.mem.sliceAsBytes(t.unconditional)); ++ // Only conditional logits are mutable. Their initial bytes live on the host, not in an immutable rig slot. ++ t.out = try t.rig.output(t.count() * 4); ++ return true; ++ } ++ ++ pub fn run(t: *Mix) !void { ++ @memcpy(t.rig.bytes(t.out, t.count() * 4), std.mem.sliceAsBytes(t.conditional)); ++ const p = t.rig.begin(); ++ p.e.setPipeline(t.pipe); ++ const Args = extern struct { count: u32, gamma: f32 }; ++ p.e.setValue(Args{ .count = @intCast(t.shape.n), .gamma = t.gamma }, 2); ++ for (0..t.shape.rows) |row| { ++ const off = row * (t.shape.n + tail) * 4; ++ p.e.setBuffer(t.out.buf, t.out.off + off, 0); ++ p.e.setBuffer(t.u.buf, t.u.off + off, 1); ++ p.e.dispatchGroups(mtl.Size.of((t.shape.n + 255) / 256, 1, 1), mtl.Size.of(256, 1, 1)); ++ } ++ try t.rig.finish(p); ++ } ++ ++ pub fn count(t: *Mix) usize { ++ return t.shape.rows * (t.shape.n + tail); ++ } ++ ++ pub fn reference(t: *Mix, want: []f64, mag: []f64) !void { ++ for (want, mag, t.conditional, t.unconditional, 0..) |*w, *m, cv, uv, i| { ++ const active = t.gamma != 1 and i % (t.shape.n + tail) < t.shape.n; ++ w.* = if (active) @as(f64, cv) + (@as(f64, t.gamma) - 1) * (@as(f64, cv) - uv) else cv; ++ m.* = if (active) @abs(cv) + (@as(f64, t.gamma) - 1) * (@abs(cv) + @abs(uv)) else 0; ++ } ++ } ++ ++ pub fn check(t: *Mix, at: g.Gate) !?[]const u8 { ++ const out = t.rig.bytes(t.out, t.count() * 4); ++ var nonfinite: usize = 0; ++ for (t.conditional, t.unconditional, 0..) |cv, uv, i| { ++ const got = c.read(.f32, out, i); ++ if (t.gamma == 1 or i % (t.shape.n + tail) >= t.shape.n) { ++ const bits = std.mem.readInt(u32, out[i * 4 ..][0..4], .little); ++ if (bits != @as(u32, @bitCast(cv))) ++ return try std.fmt.bufPrint(&t.msg, "conditional identity/codec tail changed at {d}", .{i}); ++ nonfinite += @intFromBool(!std.math.isFinite(cv)); ++ continue; ++ } ++ const alpha = @as(f64, t.gamma) - 1; ++ const want = @as(f64, cv) + alpha * (@as(f64, cv) - uv); ++ nonfinite += @intFromBool(!std.math.isFinite(want)); ++ if (c.classOf(got) != c.classOf(want)) ++ return try std.fmt.bufPrint(&t.msg, "CFG class mismatch at {d}", .{i}); ++ // Three rounded FP32 operations (or their contraction), bounded by the term magnitudes. ++ const bound = c.gamma(3, c.unit(.f32)) * (@abs(cv) + alpha * (@abs(cv) + @abs(uv))); ++ if (at != .nonfinite and std.math.isFinite(want) and @abs(got - want) > bound) ++ return try std.fmt.bufPrint(&t.msg, "CFG FP64 bound exceeded at {d}", .{i}); ++ } ++ if (at == .nonfinite and (nonfinite == 0 or nonfinite == t.count())) return "nonfinite fixture proves nothing"; ++ return null; ++ } ++}; +diff --git a/zig/tests/contract/metal_drama_combine.zig b/zig/tests/contract/metal_drama_combine.zig +new file mode 100644 +index 00000000..2e6893c0 +--- /dev/null ++++ b/zig/tests/contract/metal_drama_combine.zig +@@ -0,0 +1,285 @@ ++//! CUDA-node-compatible NVFP4 combine: independent FP64 nodes and separate pair/shared-dot sum stress. ++const std = @import("std"); ++const mtl = @import("metal"); ++const contract = @import("contract"); ++const bfr = contract.bf16_round; ++const rig_ = @import("rig.zig"); ++const source = @import("metal_drama3.zig").source; ++const c = contract.spec; ++const g = contract.gates; ++pub const Variant = enum { pairs, shared_gate }; ++fn depth(s: c.Shape) usize { ++ return (s.k + 255) / 256 + 13; ++} ++ ++pub fn spec(v: Variant) c.Spec { ++ return .{ ++ .name = switch (v) { ++ .pairs => "metal drama3 d3_nvfp4_combine pair sums", ++ .shared_gate => "metal drama3 d3_nvfp4_combine shared gate sums", ++ }, ++ .accum = .f32, ++ .depth = depth, ++ .out = .bf16, ++ .in_place = true, ++ .tiles = .{ .{}, .{}, .{ .size = 256 } }, ++ .served = &.{ .{ .rows = 8, .n = 2048, .k = 2048 }, .{ .rows = 1, .n = 2048, .k = 2048 }, .{ .rows = 168, .n = 2048, .k = 2048 } }, ++ }; ++} ++ ++const Range = struct { ++ lo: f64, ++ hi: f64, ++ fn point(v: f64) Range { ++ return .{ .lo = v, .hi = v }; ++ } ++ fn around(v: f64, e: f64) Range { ++ return .{ .lo = v - e, .hi = v + e }; ++ } ++ fn add(a: Range, b: Range) Range { ++ return .{ .lo = a.lo + b.lo, .hi = a.hi + b.hi }; ++ } ++ fn scale(a: Range, b: f64) Range { ++ return if (b >= 0) .{ .lo = a.lo * b, .hi = a.hi * b } else .{ .lo = a.hi * b, .hi = a.lo * b }; ++ } ++ fn fp32(a: Range) Range { ++ const err = c.unit(.f32) * @max(@abs(a.lo), @abs(a.hi)) + 0x1p-149; ++ return .{ .lo = a.lo - err, .hi = a.hi + err }; ++ } ++ fn bf16(a: Range) Range { ++ return .{ .lo = rounded(a.lo), .hi = rounded(a.hi) }; ++ } ++}; ++fn rounded(v: f64) f64 { ++ return bfr.bf16(@as(f32, @floatCast(v))); ++} ++fn sigmoid(v: f64) f64 { ++ return 1 / (1 + @exp(-v)); ++} ++ ++pub const Combine = struct { ++ a: std.mem.Allocator, ++ arena: std.heap.ArenaAllocator, ++ rig: rig_.Rig, ++ pipe: mtl.Pipeline, ++ variant: Variant, ++ shape: c.Shape = undefined, ++ fill: g.Fill = .random, ++ old: []u16 = &.{}, ++ shared: []u16 = &.{}, ++ x: []u16 = &.{}, ++ wg: []u16 = &.{}, ++ down: []f32 = &.{}, ++ weights: []f32 = &.{}, ++ refs: [5]rig_.Ref = undefined, ++ out: rig_.Ref = undefined, ++ ++ pub fn init(a: std.mem.Allocator, device: mtl.Device, queue: mtl.Queue, variant: Variant) !Combine { ++ const src = try source(a, "drama3 nvfp4 d3_nvfp4_combine"); ++ defer a.free(src); ++ return .{ .a = a, .arena = std.heap.ArenaAllocator.init(a), .rig = .{ .a = a, .device = device, .queue = queue }, .variant = variant, .pipe = try rig_.pipeline(device, src, "d3_nvfp4_combine") }; ++ } ++ pub fn deinit(t: *Combine) void { ++ t.arena.deinit(); ++ t.rig.deinit(); ++ t.pipe.deinit(); ++ } ++ pub fn kernel(t: *Combine) g.Kernel { ++ return .{ .spec = spec(t.variant), .ctx = t, .vt = &vt }; ++ } ++ const vt = blk: { ++ var hooks = g.hooks(Combine); ++ hooks.overlap = overlap; ++ break :blk hooks; ++ }; ++ fn overlap(_: *anyopaque) bool { ++ return true; ++ } ++ ++ pub fn prepare(t: *Combine, shape: c.Shape, fill: g.Fill, seed: u64) !bool { ++ if (shape.k == 0 or shape.k > 2048 or shape.n != shape.k) return false; ++ t.rig.reset(); ++ _ = t.arena.reset(.retain_capacity); ++ t.shape = shape; ++ t.fill = fill; ++ const a = t.arena.allocator(); ++ const n = t.count(); ++ t.old = try a.alloc(u16, n); ++ t.shared = try a.alloc(u16, n); ++ t.x = try a.alloc(u16, n); ++ t.wg = try a.alloc(u16, shape.k); ++ t.down = try a.alloc(f32, n * 8); ++ t.weights = try a.alloc(f32, shape.rows * 8); ++ var prng = std.Random.DefaultPrng.init(seed); ++ const rng = prng.random(); ++ for (t.old, t.shared, t.x) |*old, *shared, *x| { ++ old.* = c.toBf16(rng.float(f32) * 2 - 1); ++ shared.* = c.toBf16(rng.float(f32) * 2 - 1); ++ x.* = c.toBf16((rng.float(f32) * 2 - 1) / 8); ++ } ++ for (t.wg) |*v| v.* = c.toBf16((rng.float(f32) * 2 - 1) / 8); ++ for (t.down) |*v| v.* = rng.float(f32) * 8 - 4; ++ for (t.weights) |*v| v.* = (1 + rng.float(f32)) / 16; ++ if (fill == .stress) { ++ @memset(t.old, 0); ++ @memset(t.shared, c.toBf16(if (t.variant == .shared_gate) 2 else 0)); ++ @memset(t.x, 0); ++ @memset(t.wg, c.toBf16(1)); ++ @memset(t.down, 0); ++ @memset(t.weights, 0.125); ++ if (t.variant == .pairs) { ++ const terms = [_]f32{ 65536, 1, 1, 1, 1, 1, 1, -65536 }; ++ for (0..shape.rows) |row| for (terms, 0..) |v, slot| ++ @memset(t.down[(row * 8 + slot) * shape.k ..][0..shape.k], v * 8); ++ } else { ++ // Exact sum 32, magnitude 2^20+32: FP16 overflows, BF16 loses the central ones in all orders. ++ for (0..shape.rows) |row| for (0..64) |i| { ++ t.x[row * shape.k + i] = c.toBf16(if (i < 16) 32768 else if (i < 48) 1 else -32768); ++ }; ++ } ++ } ++ if (fill == .nonfinite) { ++ const values = [_]f32{ std.math.nan(f32), std.math.inf(f32), -std.math.inf(f32) }; ++ for (values, 0..) |v, row| { ++ t.down[(row * 8) * shape.k] = v; ++ t.shared[row * shape.k + 1] = c.toBf16(v); ++ t.old[row * shape.k + 2] = c.toBf16(v); ++ t.x[row * shape.k + 3] = c.toBf16(v); ++ } ++ t.x[3 * shape.k] = c.toBf16(std.math.inf(f32)); ++ t.x[3 * shape.k + 1] = c.toBf16(-std.math.inf(f32)); ++ t.wg[0] = c.toBf16(1); ++ t.wg[1] = c.toBf16(1); ++ t.wg[3] = c.toBf16(1); ++ } ++ t.out = try t.rig.output(n * 2); ++ t.refs[0] = try t.rig.input(std.mem.sliceAsBytes(t.down)); ++ t.refs[1] = try t.rig.input(std.mem.sliceAsBytes(t.weights)); ++ t.refs[2] = try t.rig.input(std.mem.sliceAsBytes(t.shared)); ++ t.refs[3] = try t.rig.input(std.mem.sliceAsBytes(t.x)); ++ t.refs[4] = try t.rig.input(std.mem.sliceAsBytes(t.wg)); ++ return true; ++ } ++ ++ pub fn run(t: *Combine) !void { ++ // Residual alone is allowed in place. Restore its original payload before each independent repeat. ++ @memcpy(t.rig.bytes(t.out, t.count() * 2), std.mem.sliceAsBytes(t.old)); ++ const p = t.rig.begin(); ++ p.e.setPipeline(t.pipe); ++ p.e.setBuffer(t.out.buf, t.out.off, 0); ++ for (t.refs, 1..) |ref, slot| p.e.setBuffer(ref.buf, ref.off, slot); ++ p.e.setValue(@as(u32, @intCast(t.shape.k)), 6); ++ p.e.dispatchGroups(mtl.Size.of(t.shape.rows, 1, 1), mtl.Size.of(256, 1, 1)); ++ try t.rig.finish(p); ++ } ++ pub fn count(t: *Combine) usize { ++ return t.shape.rows * t.shape.k; ++ } ++ ++ fn dot(t: *Combine, row: usize) struct { value: f64, magnitude: f64 } { ++ var value: f64 = 0; ++ var magnitude: f64 = 0; ++ for (t.x[row * t.shape.k ..][0..t.shape.k], t.wg) |x, w| { ++ const term = @as(f64, c.fromBf16(x)) * c.fromBf16(w); ++ value += term; ++ magnitude += @abs(term); ++ } ++ return .{ .value = value, .magnitude = magnitude }; ++ } ++ ++ fn pairs(t: *Combine, row: usize, col: usize, terms: *[8]f64) void { ++ for (terms, 0..) |*term, slot| { ++ const product = @as(f64, t.weights[row * 8 + slot]) * t.down[(row * 8 + slot) * t.shape.k + col]; ++ term.* = rounded(product); ++ } ++ } ++ fn finish(t: *Combine, row: usize, col: usize, gate: f64, sum: f64) f64 { ++ const i = row * t.shape.k + col; ++ const moe = rounded(gate * c.fromBf16(t.shared[i]) + rounded(sum)); ++ return rounded(@as(f64, c.fromBf16(t.old[i])) + moe); ++ } ++ ++ pub fn reference(t: *Combine, want: []f64, mag: []f64) !void { ++ for (0..t.shape.rows) |row| { ++ const d = t.dot(row); ++ for (0..t.shape.k) |col| { ++ var terms: [8]f64 = undefined; ++ t.pairs(row, col, &terms); ++ var sum: f64 = 0; ++ var magnitude: f64 = 0; ++ for (terms) |term| { ++ sum += term; ++ magnitude += @abs(term); ++ } ++ want[row * t.shape.k + col] = t.finish(row, col, sigmoid(d.value), sum); ++ mag[row * t.shape.k + col] = magnitude + d.magnitude; ++ } ++ } ++ } ++ ++ fn range(t: *Combine, row: usize, col: usize, gate: Range) Range { ++ var sum = Range.point(0); ++ var magnitude: f64 = 0; ++ for (0..8) |slot| { ++ const product = @as(f64, t.weights[row * 8 + slot]) * t.down[(row * 8 + slot) * t.shape.k + col]; ++ // A single FP32 product has an exact independently rounded value; CUDA rounds it again to BF16. ++ const pair = rounded(product); ++ sum = sum.add(Range.point(pair)); ++ magnitude += @abs(pair); ++ } ++ const error_sum = c.gamma(8, c.unit(.f32)) * magnitude; ++ sum = Range.around(sum.lo, error_sum).bf16(); ++ const i = row * t.shape.k + col; ++ const shared = gate.scale(c.fromBf16(t.shared[i])).fp32(); ++ const moe = shared.add(sum).fp32().bf16(); ++ return moe.add(Range.point(c.fromBf16(t.old[i]))).fp32().bf16(); ++ } ++ ++ pub fn intervals(t: *Combine, low: []f64, high: []f64) !void { ++ if (t.fill == .stress) { ++ // Integer sums/products are exact in FP32. sigmoid(32)*2 rounds to BF16 2, far from a boundary. ++ const exact: f64 = if (t.variant == .pairs) 6 else 2; ++ @memset(low, exact); ++ @memset(high, exact); ++ return; ++ } ++ for (0..t.shape.rows) |row| { ++ const d = t.dot(row); ++ const error_dot = c.gamma(@floatFromInt(depth(t.shape)), c.unit(.f32)) * d.magnitude; ++ const lo = sigmoid(d.value - error_dot); ++ const hi = sigmoid(d.value + error_dot); ++ // precise exp's two ulps, the denominator addition and reciprocal; propagate through every BF16 node. ++ const error_gate = c.gamma(6, c.unit(.f32)) * hi; ++ const gate = Range{ .lo = @max(0, lo - error_gate), .hi = @min(1, hi + error_gate) }; ++ for (0..t.shape.k) |col| { ++ const r = t.range(row, col, gate); ++ low[row * t.shape.k + col] = r.lo; ++ high[row * t.shape.k + col] = r.hi; ++ } ++ } ++ } ++ ++ pub fn bounds(t: *Combine, want: []const f64, _: []const f64, allow: []f64) !void { ++ const ranges = try t.a.alloc(f64, t.count() * 2); ++ defer t.a.free(ranges); ++ const low = ranges[0..t.count()]; ++ const high = ranges[t.count()..]; ++ try t.intervals(low, high); ++ for (allow, want, low, high) |*v, w, lo, hi| v.* = @max(@abs(w - lo), @abs(hi - w)); ++ } ++ ++ pub fn simulate(t: *Combine, acc: c.Accum, order: g.Order, out: []f64) !void { ++ const terms = try t.a.alloc(f64, t.shape.k); ++ defer t.a.free(terms); ++ for (0..t.shape.rows) |row| { ++ for (terms, t.wg, 0..) |*v, w, i| ++ v.* = @as(f64, c.fromBf16(t.x[row * t.shape.k + i])) * c.fromBf16(w); ++ const gate = sigmoid(contract.rows.sumIn(acc, order, terms)); ++ for (0..t.shape.k) |col| { ++ var pairs_: [8]f64 = undefined; ++ t.pairs(row, col, &pairs_); ++ out[row * t.shape.k + col] = t.finish(row, col, gate, contract.rows.sumIn(acc, order, &pairs_)); ++ } ++ } ++ } ++}; +diff --git a/zig/tests/contract/metal_drama_router.zig b/zig/tests/contract/metal_drama_router.zig +new file mode 100644 +index 00000000..9f964112 +--- /dev/null ++++ b/zig/tests/contract/metal_drama_router.zig +@@ -0,0 +1,224 @@ ++//! Compact routing: original-ID ordering, masking before softmax and compact slot translation. ++const std = @import("std"); ++const mtl = @import("metal"); ++const contract = @import("contract"); ++const rig_ = @import("rig.zig"); ++const source = @import("metal_drama3.zig").source; ++const c = contract.spec; ++const g = contract.gates; ++pub const Variant = enum { pruned8, pruned73, all256 }; ++const offset = 19; ++const stride = 291; ++ ++fn depth(_: c.Shape) usize { ++ return 21; ++} ++pub fn spec(v: Variant) c.Spec { ++ return .{ ++ .name = switch (v) { ++ .pruned8 => "metal drama3 d3_router_compact retained=8", ++ .pruned73 => "metal drama3 d3_router_compact retained=73", ++ .all256 => "metal drama3 d3_router_compact retained=256", ++ }, ++ .accum = .f32, ++ .depth = depth, ++ .ops = 12, ++ .out = .f32, ++ .nan = .{ .stated = "dropped nonfinites ignored; invalid softmax gives NaN weights and first8 retained original IDs" }, ++ .served = &.{ .{ .rows = 8, .n = 8, .k = 256 }, .{ .rows = 1, .n = 8, .k = 256 }, .{ .rows = 168, .n = 8, .k = 256 } }, ++ // Seven exp(-10) terms survive FP32 but both reduced-precision denominators round to 1 in every order. ++ .sums = .{ .exact = false }, ++ }; ++} ++ ++pub const Router = struct { ++ a: std.mem.Allocator, ++ rig: rig_.Rig, ++ pipe: mtl.Pipeline, ++ variant: Variant, ++ shape: c.Shape = undefined, ++ logits: []u16 = &.{}, ++ slots: [256]i32 = undefined, ++ previous: []u8 = &.{}, ++ runs: usize = 0, ++ x: rig_.Ref = undefined, ++ map: rig_.Ref = undefined, ++ picks: rig_.Ref = undefined, ++ weights: rig_.Ref = undefined, ++ msg: [200]u8 = undefined, ++ ++ pub fn init(a: std.mem.Allocator, device: mtl.Device, queue: mtl.Queue, variant: Variant) !Router { ++ const src = try source(a, "drama3 nvfp4 d3_router_compact"); ++ defer a.free(src); ++ return .{ .a = a, .rig = .{ .a = a, .device = device, .queue = queue }, .variant = variant, .pipe = try rig_.pipeline(device, src, "d3_router_compact") }; ++ } ++ pub fn deinit(t: *Router) void { ++ t.a.free(t.logits); ++ t.a.free(t.previous); ++ t.rig.deinit(); ++ t.pipe.deinit(); ++ } ++ pub fn kernel(t: *Router) g.Kernel { ++ return .{ .spec = spec(t.variant), .ctx = t, .vt = &vt }; ++ } ++ const vt = g.hooks(Router); ++ ++ pub fn prepare(t: *Router, shape: c.Shape, fill: g.Fill, seed: u64) !bool { ++ if (shape.n != 8 or shape.k != 256) return false; ++ t.rig.reset(); ++ t.shape = shape; ++ t.runs = 0; ++ t.a.free(t.logits); ++ t.logits = &.{}; ++ t.a.free(t.previous); ++ t.previous = &.{}; ++ t.logits = try t.a.alloc(u16, shape.rows * stride); ++ t.previous = try t.a.alloc(u8, t.count() * 4); ++ @memset(&t.slots, -1); ++ const kept: usize = switch (t.variant) { ++ .pruned8 => 8, ++ .pruned73 => 73, ++ .all256 => 256, ++ }; ++ // Permuted compact slots prove that ties are resolved in original ID space, not slot space. ++ for (0..kept) |slot| t.slots[(slot * 37 + 7) % 256] = @intCast(kept - 1 - slot); ++ var prng = std.Random.DefaultPrng.init(seed); ++ const rng = prng.random(); ++ @memset(t.logits, c.toBf16(91)); ++ for (0..shape.rows) |row| for (0..256) |id| { ++ const score: f32 = if (t.slots[id] < 0) 100 else if (row == 0 and fill != .stress) -2 else @as(f32, @floatFromInt(rng.uintLessThan(u32, 33))) / 8 - 2; ++ t.logits[row * stride + offset + id] = c.toBf16(score); ++ }; ++ if (fill == .stress) for (0..shape.rows) |row| { ++ var rank: usize = 0; ++ for (t.slots, 0..) |slot, id| if (slot >= 0) { ++ t.logits[row * stride + offset + id] = c.toBf16(if (rank == 0) 0 else if (rank < 8) -10 else -80); ++ rank += 1; ++ }; ++ }; ++ if (fill == .nonfinite) { ++ var ids: [256]usize = undefined; ++ var n: usize = 0; ++ for (t.slots, 0..) |slot, id| if (slot >= 0) { ++ ids[n] = id; ++ n += 1; ++ }; ++ const values = [_]f32{ std.math.nan(f32), std.math.inf(f32), -std.math.inf(f32) }; ++ for (values, 0..) |plant, row| t.logits[row * stride + offset + ids[1]] = c.toBf16(plant); ++ t.logits[3 * stride + offset + ids[0]] = c.toBf16(std.math.inf(f32)); ++ t.logits[3 * stride + offset + ids[1]] = c.toBf16(-std.math.inf(f32)); ++ if (kept < 256) for (t.slots, 0..) |slot, id| if (slot < 0) { ++ t.logits[4 * stride + offset + id] = c.toBf16(values[id % 3]); ++ }; ++ for (t.slots, 0..) |slot, id| if (slot >= 0) { ++ t.logits[5 * stride + offset + id] = c.toBf16(-std.math.inf(f32)); ++ }; ++ } ++ t.x = try t.rig.input(std.mem.sliceAsBytes(t.logits)); ++ t.map = try t.rig.input(std.mem.sliceAsBytes(&t.slots)); ++ t.picks = try t.rig.scratch(t.count() * 4); ++ t.weights = try t.rig.output(t.count() * 4); ++ return true; ++ } ++ ++ pub fn run(t: *Router) !void { ++ const p = t.rig.begin(); ++ p.e.setPipeline(t.pipe); ++ p.e.setBuffer(t.x.buf, t.x.off, 0); ++ p.e.setBuffer(t.picks.buf, t.picks.off, 1); ++ p.e.setBuffer(t.weights.buf, t.weights.off, 2); ++ p.e.setBuffer(t.map.buf, t.map.off, 3); ++ p.e.setValue([2]u32{ stride, offset }, 4); ++ p.e.dispatchGroups(mtl.Size.of(t.shape.rows, 1, 1), mtl.Size.of(256, 1, 1)); ++ try t.rig.finish(p); ++ const picked = t.rig.bytes(t.picks, t.count() * 4); ++ if (t.runs > 0 and !std.mem.eql(u8, t.previous, picked)) return error.RouterPicksNotRepeatable; ++ @memcpy(t.previous, picked); ++ t.runs += 1; ++ } ++ pub fn count(t: *Router) usize { ++ return t.shape.rows * 8; ++ } ++ ++ fn selected(t: *Router, row: usize, ids: *[8]usize) bool { ++ var used: [256]bool = @splat(false); ++ var poisoned = false; ++ var finite = false; ++ for (t.slots, 0..) |slot, id| if (slot >= 0) { ++ const v = c.fromBf16(t.logits[row * stride + offset + id]); ++ poisoned = poisoned or std.math.isNan(v) or v == std.math.inf(f32); ++ finite = finite or std.math.isFinite(v); ++ }; ++ poisoned = poisoned or !finite; ++ for (ids) |*out| { ++ var best: ?usize = null; ++ for (t.slots, 0..) |slot, id| { ++ if (slot < 0 or used[id]) continue; ++ if (best == null or (!poisoned and t.value(row, id) > t.value(row, best.?))) best = id; ++ } ++ out.* = best.?; ++ used[out.*] = true; ++ } ++ return poisoned; ++ } ++ fn value(t: *Router, row: usize, id: usize) f64 { ++ return c.fromBf16(t.logits[row * stride + offset + id]); ++ } ++ ++ pub fn reference(t: *Router, want: []f64, mag: []f64) !void { ++ for (0..t.shape.rows) |row| { ++ var ids: [8]usize = undefined; ++ const poisoned = t.selected(row, &ids); ++ var terms: [8]f64 = undefined; ++ var sum: f64 = 0; ++ for (ids, &terms) |id, *term| { ++ term.* = @exp(t.value(row, id) - t.value(row, ids[0])); ++ sum += term.*; ++ } ++ for (terms, 0..) |v, j| { ++ want[row * 8 + j] = if (poisoned) std.math.nan(f64) else v / sum; ++ mag[row * 8 + j] = @abs(want[row * 8 + j]); ++ } ++ } ++ } ++ ++ pub fn simulate(t: *Router, acc: c.Accum, order: g.Order, out: []f64) !void { ++ for (0..t.shape.rows) |row| { ++ var ids: [8]usize = undefined; ++ _ = t.selected(row, &ids); ++ var all: [256]f64 = @splat(0); ++ for (t.slots, 0..) |slot, id| if (slot >= 0) { ++ all[id] = @exp(t.value(row, id) - t.value(row, ids[0])); ++ }; ++ const total = contract.rows.sumIn(acc, order, &all); ++ var chosen: [8]f64 = undefined; ++ for (ids, &chosen) |id, *v| v.* = all[id] / total; ++ const denom = contract.rows.sumIn(acc, order, &chosen); ++ for (chosen, 0..) |v, j| out[row * 8 + j] = v / denom; ++ } ++ } ++ ++ pub fn check(t: *Router, at: g.Gate) !?[]const u8 { ++ const picked = t.rig.bytes(t.picks, t.count() * 4); ++ const weights = t.rig.bytes(t.weights, t.count() * 4); ++ const values = try t.a.alloc(f64, t.count() * 2); ++ defer t.a.free(values); ++ const want = values[0..t.count()]; ++ const mag = values[t.count()..]; ++ try t.reference(want, mag); ++ for (0..t.shape.rows) |row| { ++ var ids: [8]usize = undefined; ++ _ = t.selected(row, &ids); ++ for (ids, 0..) |id, j| { ++ const i = row * 8 + j; ++ const slot = std.mem.readInt(i32, picked[i * 4 ..][0..4], .little); ++ if (slot != t.slots[id]) return try std.fmt.bufPrint(&t.msg, "row {d} rank {d}: compact slot {d}, expected {d} (original ID {d})", .{ row, j, slot, t.slots[id], id }); ++ const got = c.read(.f32, weights, i); ++ if (c.classOf(got) != c.classOf(want[i])) return "router weight class differs from masked FP64 softmax"; ++ if (at != .nonfinite and @abs(got - want[i]) > c.bound(spec(t.variant), t.shape, want[i], mag[i])) ++ return "router weight exceeds FP64 reduction bound"; ++ } ++ } ++ return null; ++ } ++}; +diff --git a/zig/tests/contract/metal_main.zig b/zig/tests/contract/metal_main.zig +index 32eedeb6..c544e4be 100644 +--- a/zig/tests/contract/metal_main.zig ++++ b/zig/tests/contract/metal_main.zig +@@ -11,6 +11,10 @@ const limits_k = @import("metal_limits.zig"); + const h3_k = @import("metal_h3.zig"); + const h3_topk = @import("metal_h3_topk.zig"); + const drama3_k = @import("metal_drama3.zig"); ++const drama_nvfp4 = @import("metal_drama3_nvfp4.zig"); ++const drama_router = @import("metal_drama_router.zig"); ++const drama_combine = @import("metal_drama_combine.zig"); ++const drama_cfg = @import("metal_drama_cfg.zig"); + const glm_eh_k = @import("metal_glm_eh.zig"); + + const flash_pool = @import("metal_flashnext_pool.zig"); +@@ -187,6 +191,42 @@ pub fn main(init: std.process.Init) !u8 { + d3norm = try drama3_k.Norm.init(a, device, queue); + try kernels.append(a, d3norm.?.kernel()); + } ++ var d3projections: [std.enums.values(drama_nvfp4.Variant).len]drama_nvfp4.Proj = undefined; ++ var d3projection_count: usize = 0; ++ defer for (d3projections[0..d3projection_count]) |*p| p.deinit(); ++ for (std.enums.values(drama_nvfp4.Variant)) |v| { ++ if (!wanted(o, drama_nvfp4.spec(v).name)) continue; ++ d3projections[d3projection_count] = try drama_nvfp4.Proj.init(a, device, queue, v); ++ d3projection_count += 1; ++ try kernels.append(a, d3projections[d3projection_count - 1].kernel()); ++ } ++ var d3routers: [std.enums.values(drama_router.Variant).len]drama_router.Router = undefined; ++ var d3router_count: usize = 0; ++ defer for (d3routers[0..d3router_count]) |*p| p.deinit(); ++ for (std.enums.values(drama_router.Variant)) |v| { ++ if (!wanted(o, drama_router.spec(v).name)) continue; ++ d3routers[d3router_count] = try drama_router.Router.init(a, device, queue, v); ++ d3router_count += 1; ++ try kernels.append(a, d3routers[d3router_count - 1].kernel()); ++ } ++ var d3combines: [std.enums.values(drama_combine.Variant).len]drama_combine.Combine = undefined; ++ var d3combine_count: usize = 0; ++ defer for (d3combines[0..d3combine_count]) |*p| p.deinit(); ++ for (std.enums.values(drama_combine.Variant)) |v| { ++ if (!wanted(o, drama_combine.spec(v).name)) continue; ++ d3combines[d3combine_count] = try drama_combine.Combine.init(a, device, queue, v); ++ d3combine_count += 1; ++ try kernels.append(a, d3combines[d3combine_count - 1].kernel()); ++ } ++ var d3mixes: [drama_cfg.gammas.len]drama_cfg.Mix = undefined; ++ var d3mix_count: usize = 0; ++ defer for (d3mixes[0..d3mix_count]) |*p| p.deinit(); ++ for (drama_cfg.gammas) |gamma| { ++ if (!wanted(o, drama_cfg.spec(gamma).name)) continue; ++ d3mixes[d3mix_count] = try drama_cfg.Mix.init(a, device, queue, gamma); ++ d3mix_count += 1; ++ try kernels.append(a, d3mixes[d3mix_count - 1].kernel()); ++ } + var glm_eh: ?glm_eh_k.Nax = null; + defer if (glm_eh) |*n| n.deinit(); + if (wanted(o, glm_eh_k.spec().name)) { +diff --git a/zig/tests/drama_nvfp4_metal.zig b/zig/tests/drama_nvfp4_metal.zig +new file mode 100644 +index 00000000..5d5e6399 +--- /dev/null ++++ b/zig/tests/drama_nvfp4_metal.zig +@@ -0,0 +1,10 @@ ++//! Check the production NVFP4 kernels against independent references, without loading model weights. ++const std = @import("std"); ++const drama = @import("drama3"); ++ ++pub fn main() !void { ++ const gpu = try drama.gpu.Gpu.open(std.heap.page_allocator, @import("drama3_metallib").bytes); ++ defer gpu.close(); ++ try drama.quant_check.run(gpu); ++ try drama.guidance.check(gpu); ++} diff --git a/engine/nvfp4/macos/patches/0002-shared-gate-stress-regression.patch b/engine/nvfp4/macos/patches/0002-shared-gate-stress-regression.patch new file mode 100644 index 0000000..439a685 --- /dev/null +++ b/engine/nvfp4/macos/patches/0002-shared-gate-stress-regression.patch @@ -0,0 +1,124 @@ +From c59b0954129c6ee4e9c9eaa9ebc2fdea3c542f37 Mon Sep 17 00:00:00 2001 +From: jerryfane <38183696+jerryfane@users.noreply.github.com> +Date: Sun, 25 Oct 2026 05:16:43 +0400 +Subject: [PATCH] tests: discriminate shared-gate sums after sigmoid + +Center the exact FP32 dot at one so sigmoid saturation cannot hide a +lower-precision accumulator's overflow. Production kernels and tolerance +widths are unchanged. + +The independent CPU regression passed for FP16 and BF16 in all three sum +orders. Only the corrected failed GPU variant was rerun on Apple M5 Ultra: +1/1 passed, retaining the other 13 variants and runtime-limit receipts. + +--- + zig/build/kernel_contract.zig | 7 ++++-- + zig/tests/contract/drama3_nvfp4_ref.zig | 25 ++++++++++++++++++++++ + zig/tests/contract/metal_drama_combine.zig | 13 +++++------ + 3 files changed, 37 insertions(+), 8 deletions(-) + +diff --git a/zig/build/kernel_contract.zig b/zig/build/kernel_contract.zig +index fcac3fff..d8a7ca1c 100644 +--- a/zig/build/kernel_contract.zig ++++ b/zig/build/kernel_contract.zig +@@ -36,14 +36,17 @@ pub fn hostTests(b: *std.Build, family: *std.Build.Step) void { + }) })); + b.step("test-delta-slots-contract-host", "FP64 DeltaNet and slot corruption controls without shaders").dependOn(&delta.step); + step.dependOn(&delta.step); +- const nvfp4 = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ ++ const nvfp4_module = b.createModule(.{ + .root_source_file = b.path("zig/tests/contract/drama3_nvfp4_ref.zig"), + .target = b.graph.host, + .optimize = .debug, + .imports = &.{.{ .name = "contract", .module = module(b, b.graph.host, .debug) }}, +- }) })); ++ }); ++ const nvfp4 = b.addRunArtifact(b.addTest(.{ .root_module = nvfp4_module })); + b.step("test-drama-nvfp4-reference-host", "NVFP4 scalar reference and reduced-accumulator discrimination").dependOn(&nvfp4.step); + step.dependOn(&nvfp4.step); ++ const shared_gate = b.addRunArtifact(b.addTest(.{ .root_module = nvfp4_module, .filters = &.{"shared gate stress"} })); ++ b.step("test-drama-combine-stress-host", "Shared-gate stress remains discriminating after nonlinear output").dependOn(&shared_gate.step); + } + + /// macOS: tf-kernel-contract over the engine's Metal kernels; `kernel-contract` runs it on this Mac's GPU. +diff --git a/zig/tests/contract/drama3_nvfp4_ref.zig b/zig/tests/contract/drama3_nvfp4_ref.zig +index 6ee80a54..ac766bc6 100644 +--- a/zig/tests/contract/drama3_nvfp4_ref.zig ++++ b/zig/tests/contract/drama3_nvfp4_ref.zig +@@ -182,6 +182,31 @@ pub fn upRange(gate: Interval, up: Interval) Interval { + return fp32(g.mul(inverse)).mul(u).rounded(); + } + ++pub const shared_gate_terms = 65; ++ ++pub fn sharedGateStressTerm(i: usize) f32 { ++ return if (i < 16) 32768 else if (i < 48) 1 else if (i < 64) -32768 else if (i == 64) -31 else 0; ++} ++ ++pub fn sharedGateOutput(dot: f64) f64 { ++ return bf(2 / (1 + @exp(-dot))); ++} ++ ++test "shared gate stress remains discriminating after sigmoid and BF16 output" { ++ var terms: [2048]f64 = @splat(0); ++ for (terms[0..shared_gate_terms], 0..) |*v, i| v.* = sharedGateStressTerm(i); ++ const want = sharedGateOutput(1); ++ try std.testing.expectEqual(@as(f64, 1.4609375), want); ++ for ([_]rows.Order{ .sequential, .pairwise, .strided }) |order| { ++ try std.testing.expectEqual(@as(f64, 1), rows.sumIn(.f32, order, &terms)); ++ for ([_]c.Accum{ .f16, .bf16 }) |acc| { ++ const got = sharedGateOutput(rows.sumIn(acc, order, &terms)); ++ // Exact FP32 sums yield one stable BF16 value, so the stress interval has zero width. ++ try std.testing.expect(!std.math.isFinite(got) or got != want); ++ } ++ } ++} ++ + test "NVFP4 independent decoding includes signs, subnormal scales, signed zero and maximum finite scale" { + try std.testing.expectEqual(@as(f64, 0.5), fp4(1)); + try std.testing.expectEqual(@as(f64, -6), fp4(15)); +diff --git a/zig/tests/contract/metal_drama_combine.zig b/zig/tests/contract/metal_drama_combine.zig +index 2e6893c0..21291a6b 100644 +--- a/zig/tests/contract/metal_drama_combine.zig ++++ b/zig/tests/contract/metal_drama_combine.zig +@@ -5,6 +5,7 @@ const contract = @import("contract"); + const bfr = contract.bf16_round; + const rig_ = @import("rig.zig"); + const source = @import("metal_drama3.zig").source; ++const ref = @import("drama3_nvfp4_ref.zig"); + const c = contract.spec; + const g = contract.gates; + pub const Variant = enum { pairs, shared_gate }; +@@ -132,9 +133,9 @@ pub const Combine = struct { + for (0..shape.rows) |row| for (terms, 0..) |v, slot| + @memset(t.down[(row * 8 + slot) * shape.k ..][0..shape.k], v * 8); + } else { +- // Exact sum 32, magnitude 2^20+32: FP16 overflows, BF16 loses the central ones in all orders. +- for (0..shape.rows) |row| for (0..64) |i| { +- t.x[row * shape.k + i] = c.toBf16(if (i < 16) 32768 else if (i < 48) 1 else -32768); ++ // Exact dot 1 avoids sigmoid saturation hiding a reduced-precision sum's overflow. ++ for (0..shape.rows) |row| for (0..ref.shared_gate_terms) |i| { ++ t.x[row * shape.k + i] = c.toBf16(ref.sharedGateStressTerm(i)); + }; + } + } +@@ -167,7 +168,7 @@ pub const Combine = struct { + const p = t.rig.begin(); + p.e.setPipeline(t.pipe); + p.e.setBuffer(t.out.buf, t.out.off, 0); +- for (t.refs, 1..) |ref, slot| p.e.setBuffer(ref.buf, ref.off, slot); ++ for (t.refs, 1..) |buffer, slot| p.e.setBuffer(buffer.buf, buffer.off, slot); + p.e.setValue(@as(u32, @intCast(t.shape.k)), 6); + p.e.dispatchGroups(mtl.Size.of(t.shape.rows, 1, 1), mtl.Size.of(256, 1, 1)); + try t.rig.finish(p); +@@ -237,8 +238,8 @@ pub const Combine = struct { + + pub fn intervals(t: *Combine, low: []f64, high: []f64) !void { + if (t.fill == .stress) { +- // Integer sums/products are exact in FP32. sigmoid(32)*2 rounds to BF16 2, far from a boundary. +- const exact: f64 = if (t.variant == .pairs) 6 else 2; ++ // Exact FP32 integer sums stay away from BF16 rounding boundaries after the sigmoid. ++ const exact: f64 = if (t.variant == .pairs) 6 else ref.sharedGateOutput(1); + @memset(low, exact); + @memset(high, exact); + return; diff --git a/engine/nvfp4/macos/patches/0003-shared-cli-extension-point.patch b/engine/nvfp4/macos/patches/0003-shared-cli-extension-point.patch new file mode 100644 index 0000000..40bfa87 --- /dev/null +++ b/engine/nvfp4/macos/patches/0003-shared-cli-extension-point.patch @@ -0,0 +1,35 @@ +From 6585530054e273aa1a1554549585489f2497cb41 Mon Sep 17 00:00:00 2001 +From: jerryfane <38183696+jerryfane@users.noreply.github.com> +Date: Sun, 25 Oct 2026 05:47:21 +0400 +Subject: [PATCH] drama: keep shared flag-table extension points free + +Group the independent CFG flags beside voice selection, leaving both common table ends available to pending native Notes additions. Flag names, values and parsing policy are unchanged; only capability-list order changes. No kernel, arithmetic or tolerance changes. +--- + zig/src/server/cli.zig | 4 +++- + 1 file changed, 3 insertions(+), 1 deletion(-) + +diff --git a/zig/src/server/cli.zig b/zig/src/server/cli.zig +index 1aa949e2..94b9ed37 100644 +--- a/zig/src/server/cli.zig ++++ b/zig/src/server/cli.zig +@@ -49,6 +49,7 @@ pub const flags = [_]Flag{ + .{ .name = "--chat-template", .native = true }, + .{ .name = "--voices", .native = true }, + .{ .name = "--default-voice", .native = true }, ++} ++ @import("drama_cfg.zig").flags ++ [_]Flag{ + .{ .name = "--alias", .kind = .append, .native = true }, + .{ .name = "--api-key", .kind = .append, .native = true }, + .{ .name = "--api-key-file", .native = true }, +@@ -148,7 +149,8 @@ pub const flags = [_]Flag{ + .{ .name = "--recall-read-only", .kind = .store_true, .native = true }, + // the native server's own: Cortheon's reasoning program and server-run web research for every chat request + .{ .name = "--cortheon", .kind = .store_true, .native = true }, +-} ++ @import("drama_cfg.zig").flags; ++}; ++ + /// The variables this binary honours as the Python engine does, then the CUDA build's. + pub const env = [_][]const u8{ "TENSORFOLD_API_KEY", "TENSORFOLD_NO_LIVE", "TENSORFOLD_SEED_SALT", "TENSORFOLD_REQUEST_LOG", "TENSORFOLD_NO_UPDATE_CHECK", "HF_HOME", "HF_HUB_CACHE", "HF_HUB_OFFLINE", "BRAVE_SEARCH_API_KEY" } ++ + (if (cuda_build) [_][]const u8{ "TF_CUDA_DEVICE", "TF_CUDA_SEGMENTS", "TENSORFOLD_CUDA_KERNELS", "TENSORFOLD_MEMORY_RESERVE_GIB", "TENSORFOLD_CUDA_MEMORY_LIMIT_GB" } else [_][]const u8{}); +-- +2.54.0 (Apple Git-157) + diff --git a/engine/nvfp4/macos/patches/SHA256SUMS b/engine/nvfp4/macos/patches/SHA256SUMS new file mode 100644 index 0000000..22224d6 --- /dev/null +++ b/engine/nvfp4/macos/patches/SHA256SUMS @@ -0,0 +1,3 @@ +f282318f0e8dae43b19d1fa93ec05691696a0352eaeac41f9ffc334690d8b388 0001-drama3-metal-nvfp4-cfg.patch +0a68a3b66da1991f1a04980cb6958e4acc8856a88efdbe1fe10adc6df588cfc8 0002-shared-gate-stress-regression.patch +794a5b6008487ec705cadb098eb1c15737da743aee0855b610510506a0ce38db 0003-shared-cli-extension-point.patch diff --git a/engine/patches/README.md b/engine/patches/README.md index 7e7b88d..36e20a4 100644 --- a/engine/patches/README.md +++ b/engine/patches/README.md @@ -17,3 +17,18 @@ boundary including tower lookahead. It preserves uncached chunk cuts and CUDA ex Its CPU regressions cover identity, cancellation/error invalidation, allocation ownership and prompt boundaries; real-model transition/corpus/stream checks are described in [`recipe/README-macos.md`](../../recipe/README-macos.md#correctness-gates). + +## Native macOS NVFP4 series + +On Darwin, `--weights nvfp4` uses the separate pin in +[`../nvfp4/macos/UPSTREAM`](../nvfp4/macos/UPSTREAM) and the ordered, checksummed series below. +It does not change the full-model series above or the CUDA NVFP4 pin. + +| Patch | Scope | Upstream | +|---|---|---| +| [`0001-drama3-metal-nvfp4-cfg.patch`](../nvfp4/macos/patches/0001-drama3-metal-nvfp4-cfg.patch) | Packed NVFP4 experts, compact routing, native reference CFG and registered numerical contracts | [TensorFold #1268](https://github.com/ashhart/TensorFold/pull/1268) | +| [`0002-shared-gate-stress-regression.patch`](../nvfp4/macos/patches/0002-shared-gate-stress-regression.patch) | Test-only correction that distinguishes reduced-precision accumulation after sigmoid; no kernel or tolerance change | [TensorFold #1268](https://github.com/ashhart/TensorFold/pull/1268) | +| [`0003-shared-cli-extension-point.patch`](../nvfp4/macos/patches/0003-shared-cli-extension-point.patch) | Group CFG flags beside voice selection without occupying shared CLI-table extension points; no parsing or numerical change | [TensorFold #1268](https://github.com/ashhart/TensorFold/pull/1268) | + +The [Mac qualification receipt](../../results/macos-nvfp4.json) separates observed GPU/model results, +the initial failed fixture and focused correction, and the final CPU-only CLI integration build. diff --git a/recipe/README-macos.md b/recipe/README-macos.md index ef7f5cb..fa28724 100644 --- a/recipe/README-macos.md +++ b/recipe/README-macos.md @@ -1,9 +1,9 @@ # macOS: running the pinned recipe on Apple silicon -The full-model Metal build uses the same pinned TensorFold commit, ordered patch series, -checkpoint revision and HTTP contract as the Linux recipe. It does not require CUDA, -`nvcc`, systemd, or `/proc`. Use `recipe/pilot_macos.py`; the Linux launcher remains unchanged. -The separate NVFP4 engine/weights are not qualified by this Mac guide. +Both `--weights full` and `--weights nvfp4` use the native Zig/Metal engine, not Python/MLX. +The full model retains the Linux recipe's existing pin and patches. NVFP4 uses the separate +[Mac engine pin](../engine/nvfp4/macos/UPSTREAM) and patch series; the CUDA NVFP4 pin is unchanged. +Neither path requires CUDA, `nvcc`, systemd, or `/proc`. Use `recipe/pilot_macos.py`. This is a local, cooperative pilot—not a production service installer, GPU reservation system, memory sandbox, or proof that the whole host is idle. Do not interrupt another @@ -15,23 +15,45 @@ application to make an admission check pass. - Zig **0.17.0**. The official `zig-aarch64-macos-0.17.0.tar.xz` archive has SHA-256 `b607e9b9234790a008116ae5bdb71c6243b84b9fb42a53a9e70fde41c06c536a`. - Python 3.9 or newer for the stdlib-only launcher and validation tools. -- Approximately 39 GiB of checkpoint storage plus build/cache space, and sufficient - free + inactive memory for admission. Disk size is not a resident-memory prediction. +- The Hugging Face CLI (`hf`) on `PATH`. Install it in an isolated Python environment before downloading: + + ```sh + python3 -m venv .venv-hf + . .venv-hf/bin/activate + python -m pip install huggingface_hub + ``` + +- Approximately 21 GiB for the NVFP4 checkpoint, or 39 GiB for the full checkpoint, plus + build/cache space and sufficient free + inactive memory. Disk size is not a resident-memory prediction. - Permission to use the checkpoint under the [pinned model license](https://huggingface.co/fishaudio/drama-3/blob/e0a9218aa48888270f06ef1ed9ea67622c0e17d2/LICENSE). -From the repository root: +From the repository root, for the NVFP4 checkpoint: + +```sh +hf download plotarmordev/drama-3-NVFP4-Pruned \ + --revision 494db08481db952f834b2241bce63fb3dfa72554 --local-dir drama-3-NVFP4-Pruned +(cd drama-3-NVFP4-Pruned && shasum -a 256 -c ../recipe/checksums/nvfp4.sha256) +ZIG=/absolute/path/to/zig-0.17.0/zig JOBS=2 scripts/build-engine.sh --weights nvfp4 +``` + +All fifteen checkpoint files must report `OK`. The launcher also verifies the selected manifest +on every start, including an explicit `--checkpoint` override; an incomplete or changed directory +is refused rather than repaired. An absent default directory can be downloaded automatically. +The build verifies its ordered patch checksums and does not require a CUDA compiler on Darwin. +Outputs are `build/tensorfold-nvfp4/zig-out/native/bin/tensorfold-native` and +`build/tensorfold-nvfp4/zig-out/bin/{drama3-generate,drama3-voice}`. + +For the unchanged full-model recipe: ```sh hf download fishaudio/drama-3 --revision e0a9218aa48888270f06ef1ed9ea67622c0e17d2 --local-dir drama-3 (cd drama-3 && shasum -a 256 -c ../recipe/checksums/checkpoint.sha256) -ZIG=/absolute/path/to/zig-0.17.0/zig scripts/build-engine.sh --weights full +ZIG=/absolute/path/to/zig-0.17.0/zig JOBS=2 scripts/build-engine.sh --weights full ``` -All ten checkpoint files must report `OK` before a model run. The build script checks the -ordered patches and omits the Linux/Windows-only `-Dnvcc` option on Darwin. Outputs: - -- `build/tensorfold/zig-out/native/bin/tensorfold-native` -- `build/tensorfold/zig-out/bin/drama3-generate` +All ten full-model files must report `OK`. Its server and `drama3-generate` remain under +`build/tensorfold/zig-out/`. A build directory belonging to another engine repository is refused: +pass a fresh positional `BUILD_DIR`, then select its server explicitly with `--native-bin`. Use the installed build's `serve --help` for its supported options; do not assume features from newer upstream releases exist in this pinned tree. @@ -45,16 +67,28 @@ unshared local pilot, the following creates its local runtime directory and lock ```sh mkdir -p runtime touch runtime/gpu.lock -python3 recipe/pilot_macos.py \ - --native-bin build/tensorfold/zig-out/native/bin/tensorfold-native \ - --checkpoint drama-3 --runtime-dir runtime --gpu-lock runtime/gpu.lock \ - --evidence-dir runs/macos-pilot-01 +python3 recipe/pilot_macos.py --weights nvfp4 \ + --runtime-dir runtime --gpu-lock runtime/gpu.lock --evidence-dir runs/macos-pilot-01 ``` The default bind address is **127.0.0.1**, port **18793**; `--host` accepts numeric IPv4 addresses. An occupied port or lock is a refusal, not permission to displace its owner. This unauthenticated API should not be exposed casually; LAN placement does not exempt a hosted service from the model license. +Use `--weights full` for the full model. Explicit `--native-bin` and `--checkpoint` remain available; +with `--weights` the selected checkpoint manifest is still mandatory. Without `--weights`, the legacy +explicit-binary/checkpoint mode makes no pinned-manifest verification claim. `session.json` records +the selection, checkpoint revision, manifest, actual binary and command. + +NVFP4 defaults to `--cfg-gamma 1.5 --cfg-batched`. On Metal, paired guidance means two independent +backbone states submitted in one command buffer, sharing weights; it is not a two-row/vectorized GEMM. +Only semantic logits are guided; residual codec heads stay conditional. Builtin voices remain unguided. +Admission conservatively budgets the additional guidance state. To replace these defaults, pass +`-- --cfg-gamma 1.5` for separate submissions or `-- --cfg-gamma 0` to disable guidance. +An explicit `--cfg-gamma`, `--cfg-batched`, or `--cfg-allow-request-gamma` replaces the default group. +Finite gamma values 1–5 are supported, with 1 preserving conditional logits; 0 disables the policy. +Request `X-Cfg-Gamma` overrides require `--cfg-allow-request-gamma`. `--cfg-every` is unsupported. + The launcher records `session.json`, `server.log`, `memory.jsonl`, `ready.json` and `completion.json`. `DRAMA_PILOT_READY` means the launched server passed its health check, not that audio correctness or performance has been verified. Explicit `--evidence-reuse` clears stale @@ -105,7 +139,23 @@ Use a synthetic voice, your own voice, or a voice for which you have the require The caller must attest the source explicitly; the catalog writer does not invent consent. Reference audio, codec rows, and generated voice catalogs stay out of Git. -Encode a matching transcript and a clean single-speaker reference clip: +For the NVFP4 Mac engine, use its native attested catalog writer: + +```sh +mkdir -p runtime/voices +build/tensorfold-nvfp4/zig-out/bin/drama3-voice \ + --codec drama-3-NVFP4-Pruned/codec.pth --ref-audio clip.wav --ref-text "$(cat clip.txt)" \ + --source-kind self --source-note "My own recording; permission retained locally." \ + --out runtime/voices/my-voice.json \ + --generate "$PWD/build/tensorfold-nvfp4/zig-out/bin/drama3-generate" +``` + +`--ref-text` takes literal text, so the command reads the exact transcript from `clip.txt`. Encoding +is GPU work and needs a coordinated window with the speech server stopped. The writer emits +`drama3-http-voice/1`; the older CUDA quant engine's `tools/zig/drama_voice.py` is not a Mac encoder. +The Mac speech API accepts attested catalog voices, not request-supplied reference audio or codes. + +For the full-model engine, encode the reference and write the same attested catalog format: ```sh build/tensorfold/zig-out/bin/drama3-generate --codec drama-3/codec.pth \ @@ -130,11 +180,13 @@ Start a new coordinated pilot with `-- --voices /absolute/path/to/runtime/voices directory. `GET /v1/voices` lists loaded voices. Without `--default-voice`, omitted/default requests use the builtin voice. Setting `--default-voice NAME` changes both cases. -Patch 0012 reuses one reference-prefix snapshot per Metal lane. A different or absent prefix -invalidates the live cache before prefill overwrites it; switching voices reuses the allocated -snapshot buffers rather than retaining a new allocation per switch. Prefix identity includes -the prompt boundary and the audio tower's lookahead dependency. Generation settings and -sampling are unchanged. This is an output-preserving optimization, not a multi-voice cache. +The full-model patch 0012 and the NVFP4 Mac base reuse one conditional reference-prefix snapshot per +Metal lane. A different or absent prefix invalidates it before prefill; switching voices reuses the +allocated snapshot buffers. Guided requests reset their independent unconditional backbone state; +cancellation invalidates conditional retention before another request. Neither is a multi-voice cache. +The NVFP4 checkpoint stays packed: E2M1 expert codes, block-16 E4M3 scales and F32 global scales, +with BF16 dense/tower weights and FP32 accumulation. No checkpoint conversion or BF16 expert expansion +is performed. Pruned routing masks dropped original IDs before softmax and retains original-ID tie order. ## Correctness gates @@ -185,3 +237,30 @@ reference-boundary variants and two-lane parity passed. Eight warmed voice switc 0 KiB process-RSS spread on each build. RSS is not direct Metal allocation telemetry. All three servers and their owned helpers settled. This correctness evidence makes no speed, CUDA-runtime or NVFP4-Mac claim. + +### Recorded NVFP4 qualification + +[`results/macos-nvfp4.json`](../results/macos-nvfp4.json) records the exact NVFP4 revision above on +an Apple M5 Ultra with native TensorFold 1.1.4. The public build entry point produced the server, +codec generator and voice writer from the pinned base and verified patch. A second, test-only +patch corrects an accumulation stress fixture; production kernels and tolerances are unchanged. +A final CLI-only patch groups CFG flags beside voice selection to avoid an upstream integration +collision. Its native build, version and help were exercised separately; the completed GPU checks +were not replayed. The receipt distinguishes both binary hashes and their verification scopes. + +The initial Metal contract run passed 13 of 14 Drama variants and all runtime resource limits. +Its shared-gate fixture failed to distinguish FP16 accumulation after sigmoid saturation. +An independent CPU regression now proves that distinction; only the corrected failed variant +was rerun and passed. The receipt retains both results, not a fabricated full-suite pass. + +Three sequential single-stream servers completed 28 HTTP checks: 15 complete PCM responses, +one early stream cancellation and 12 expected validation refusals. Paired and separate gamma-1.5 +submissions produced identical PCM for builtin and two synthetic reference voices. Repetition, +reference transitions, streamed/whole identity, the post-cancellation response, and opt-in gamma-1/2 +repeatability passed. All original servers exited zero after their planned stop, and the supervising +process groups settled. No installed runtime or production service was changed. + +These are correctness checks, not a speed or voice-quality comparison. Live qualification used +the native engine directly under the existing cooperative guard; the launcher's selected-weight +path has CPU coverage. M1–M4 execution, smaller-memory Macs, CUDA parity, the full multi-family +Metal suite and offline M1/M2 compiler receipts remain unqualified. diff --git a/recipe/README.md b/recipe/README.md index b41c0a6..e0aaf6b 100644 --- a/recipe/README.md +++ b/recipe/README.md @@ -3,6 +3,10 @@ `pilot.py` starts one Drama 3 server, native or reference, as a transient `systemd --user` unit, and stops it cleanly. It needs Linux, user systemd and Python 3 (standard library only). +On Apple silicon use [`pilot_macos.py`](README-macos.md), with the same explicit `--weights` +selection and checkpoint manifests. It uses native Metal and owned process groups, not systemd. +Its NVFP4 engine pin is separate from CUDA; the full-model pin remains unchanged. + ```sh python3 recipe/pilot.py --weights nvfp4 \ --runtime-dir "$RUNTIME_DIR" \ @@ -17,7 +21,10 @@ revision in `sources.json`, into `drama-3-NVFP4-Pruned/` or `drama-3/` at the re hashes every pinned file and refuses, before it takes the GPU lock, if one is missing or changed. With `nvfp4` it adds `--cfg-gamma 1.5 --cfg-batched`, the guidance the quant was measured with, unless a `--cfg-gamma`, `--cfg-allow-request-gamma` or `--cfg-batched` of your own follows `--`. `--checkpoint` and `--native-bin` override -the directory and the executable; without `--weights`, pass both, and nothing is downloaded or added. +the directory and executable, but a selected `--weights` manifest is still checked against an overridden +checkpoint. Only an absent default checkpoint is downloaded; explicit paths and existing incomplete directories +are refused rather than repaired. Without `--weights`, pass both paths; nothing is downloaded or added, +and no pinned-checkpoint verification is claimed. For the SGLang-Omni reference, replace `--native-bin …` with `--backend sglang-omni`. The reference also needs `venv/` (with `sgl-omni` installed) and an `sglang-omni/` checkout at the revision in `sources.json`, both inside diff --git a/recipe/pilot.py b/recipe/pilot.py index d817481..ceb0a04 100644 --- a/recipe/pilot.py +++ b/recipe/pilot.py @@ -4,11 +4,9 @@ import argparse import fcntl -import hashlib import json import os from pathlib import Path -import shutil import signal import socket import subprocess @@ -17,56 +15,13 @@ import urllib.error import urllib.request +import weight_selection + GIB = 1024**3 HEADROOM = 30 * GIB STOP_MARGIN = 4 * GIB CGROUP_CAP = 54 * GIB CUDA_HOME = Path(os.environ.get('CUDA_HOME', '/usr/local/cuda')) -RECIPE = Path(__file__).resolve().parent -REPO = RECIPE.parent -# What --weights chooses: the pinned checkpoint (a sources.json entry) and its local directory and pins, the engine -# build scripts/build-engine.sh makes for it, and the serve flags it was measured with (results/full-vs-quant-01). -WEIGHTS = { - 'full': {'source': 'checkpoint', 'dir': 'drama-3', 'pins': 'checkpoint.sha256', 'build': 'tensorfold', - 'serve': ()}, - 'nvfp4': {'source': 'nvfp4_checkpoint', 'dir': 'drama-3-NVFP4-Pruned', 'pins': 'nvfp4.sha256', - 'build': 'tensorfold-nvfp4', 'serve': ('--cfg-gamma', '1.5', '--cfg-batched')}, -} -# A serve flag after -- that sets guidance replaces the measured guidance flags as a whole. -GUIDANCE_FLAGS = ('--cfg-gamma', '--cfg-allow-request-gamma', '--cfg-batched') - - -def sha256(path: Path) -> str: - digest = hashlib.sha256() - with path.open('rb') as file: - for block in iter(lambda: file.read(1 << 24), b''): - digest.update(block) - return digest.hexdigest() - - -def fetch_checkpoint(choice: str) -> Path: - """The pinned checkpoint for --weights CHOICE: downloaded on first use, checked against its pins on every start.""" - spec = WEIGHTS[choice] - source = json.loads((RECIPE / 'sources.json').read_text())[spec['source']] - target = REPO / spec['dir'] - pins = RECIPE / 'checksums' / spec['pins'] - entries = [line.split(None, 1) for line in pins.read_text().splitlines() if line.strip()] - if not target.exists(): - hf = shutil.which('hf') - if hf is None: - raise SystemExit('pilot: the hf command is missing; install it with: python3 -m pip install -U huggingface_hub') - print(f'Downloading {source["repository"]} at {source["revision"][:8]} into {target}', flush=True) - subprocess.run([hf, 'download', source['repository'], '--revision', source['revision'], - '--local-dir', str(target)], check=True) - print(f'Checking {target} against {pins.relative_to(REPO)}', flush=True) - started = time.monotonic() - bad = [name for digest, name in entries if not (target / name).is_file() or sha256(target / name) != digest] - if bad: - raise SystemExit(f'pilot: {target} does not match {pins}: {", ".join(bad)} missing or changed; no model was ' - f'started. Repair it with: hf download {source["repository"]} --revision ' - f'{source["revision"]} --local-dir {target}, or delete it and run again') - print(f'{target} matches {pins.relative_to(REPO)} ({time.monotonic() - started:.0f} s)', flush=True) - return target def memory_available() -> int: @@ -78,7 +33,7 @@ def memory_available() -> int: def main() -> int: parser = argparse.ArgumentParser(description=__doc__) - parser.add_argument('--weights', choices=tuple(WEIGHTS), + parser.add_argument('--weights', choices=tuple(weight_selection.WEIGHTS), help='Download (first run), verify and serve these pinned weights with the engine ' 'scripts/build-engine.sh --weights builds for them; nvfp4 also applies its measured ' 'guidance, --cfg-gamma 1.5 --cfg-batched, unless a guidance flag follows --') @@ -90,7 +45,7 @@ def main() -> int: parser.add_argument('--runtime-dir', type=Path, required=True, help='Isolated runtime and cache directory') parser.add_argument('--checkpoint', type=Path, - help='Checkpoint directory, verified beforehand (default with --weights: its download)') + help='Checkpoint directory (with --weights: verified against the selected pins)') parser.add_argument('--gpu-lock', type=Path, required=True, help='Existing lock file that every GPU workload on this host takes; never a second lock') parser.add_argument('--evidence-dir', type=Path, required=True) @@ -106,16 +61,13 @@ def main() -> int: if args.weights == 'nvfp4' and args.backend != 'native': parser.error('Only the native backend serves the NVFP4 weights') root = args.runtime_dir.resolve() - checkpoint = (args.checkpoint or REPO / WEIGHTS[args.weights]['dir']).resolve() + checkpoint = weight_selection.checkpoint_path(args.weights, args.checkpoint) evidence = args.evidence_dir.resolve() if args.backend == 'native': if args.native_bin is None and args.weights is None: parser.error('The native backend requires --native-bin, or --weights') - native_bin = args.native_bin or REPO / 'build' / WEIGHTS[args.weights]['build'] / 'zig-out/native/bin/tensorfold-native' - executable = native_bin.resolve() - serve_args = list(args.serve_args) - if args.weights and not any(a.split('=', 1)[0] in GUIDANCE_FLAGS for a in serve_args): - serve_args[:0] = WEIGHTS[args.weights]['serve'] + executable = weight_selection.native_path(args.weights, args.native_bin) + serve_args = weight_selection.serve_arguments(args.weights, args.serve_args) working_directory = root server_command = [str(executable), 'serve', str(checkpoint), '--host', '127.0.0.1', '--port', str(args.port), '--name', 'fishaudio/drama-3', @@ -147,8 +99,10 @@ def main() -> int: backend_environment.append('LD_LIBRARY_PATH=' + ':'.join(map(str, directories))) if not args.gpu_lock.is_file(): parser.error('The operator-specified shared GPU lock must already exist') - if args.checkpoint is None: - fetch_checkpoint(args.weights) + try: + provenance = weight_selection.prepare_checkpoint(args.weights, args.checkpoint) + except (OSError, ValueError, subprocess.CalledProcessError) as exc: + parser.error(str(exc)) evidence.mkdir(parents=True, exist_ok=False) weights = sum(p.stat().st_size for p in checkpoint.glob('*.safetensors')) weights += (checkpoint / 'codec.pth').stat().st_size @@ -194,7 +148,7 @@ def main() -> int: f'FLASHINFER_WORKSPACE_BASE={root / "runtime-cache/flashinfer"}', *backend_environment, *server_command] - initial = {'backend': args.backend, 'weights_bytes': weights, 'initial_memavailable_bytes': available, + initial = {**provenance, 'backend': args.backend, 'weights_bytes': weights, 'initial_memavailable_bytes': available, 'required_memavailable_bytes': required, 'cgroup_memory_max_bytes': CGROUP_CAP, 'headroom_floor_bytes': HEADROOM, 'abort_threshold_bytes': HEADROOM + STOP_MARGIN, 'unit': unit, 'server_command': server_command, 'launch_command': command, diff --git a/recipe/pilot_macos.py b/recipe/pilot_macos.py index 27baf68..9ffe160 100644 --- a/recipe/pilot_macos.py +++ b/recipe/pilot_macos.py @@ -25,6 +25,8 @@ import subprocess import time +import weight_selection + GIB = 1024**3 HEADROOM = 30 * GIB STOP_MARGIN = 4 * GIB @@ -132,10 +134,15 @@ def main() -> int: parser = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) parser.add_argument('--backend', choices=('native', 'sglang-omni'), default='native') - parser.add_argument('--native-bin', type=Path, help='Patched tensorfold-native executable') + parser.add_argument('--weights', choices=tuple(weight_selection.WEIGHTS), + help='Download (first run), verify and serve pinned weights; nvfp4 adds ' + '--cfg-gamma 1.5 --cfg-batched unless guidance follows --') + parser.add_argument('--native-bin', type=Path, + help='Patched executable (default with --weights: its engine in build/)') parser.add_argument('--runtime-dir', type=Path, required=True, help='Isolated runtime and cache directory') - parser.add_argument('--checkpoint', type=Path, required=True) + parser.add_argument('--checkpoint', type=Path, + help='Checkpoint directory (with --weights: verified against the selected pins)') parser.add_argument('--gpu-lock', type=Path, required=True, help='Existing lock file that every GPU workload on this host takes') parser.add_argument('--evidence-dir', type=Path, required=True) @@ -154,9 +161,11 @@ def main() -> int: args = parser.parse_args() if not args.gpu_lock.is_file(): parser.error('The operator-specified shared GPU lock must already exist') + if args.weights is None and args.checkpoint is None: + parser.error('Pass --weights nvfp4 or --weights full, or --checkpoint DIR') root = args.runtime_dir.resolve() - checkpoint = args.checkpoint.resolve() + checkpoint = weight_selection.checkpoint_path(args.weights, args.checkpoint) evidence = args.evidence_dir.resolve() with args.gpu_lock.open('r+') as lock: @@ -212,16 +221,15 @@ def request_stop(signum: int, _frame: object) -> None: socket.inet_pton(socket.AF_INET, args.host) if args.backend != 'native': raise ValueError('sglang-omni needs the CUDA reference stack; this launcher is native-only') - if args.native_bin is None: - raise ValueError('The native backend requires --native-bin') if args.min_available_gib is not None and args.min_available_gib < ABORT_AVAILABLE / GIB: raise ValueError('Admission override cannot be below the 34 GiB abort threshold') for value in (args.reserve_gib, args.min_available_gib): if value is not None and not math.isfinite(value * GIB): raise ValueError('GiB argument is too large') - executable = args.native_bin.resolve() + executable = weight_selection.native_path(args.weights, args.native_bin) if not executable.is_file() or not os.access(executable, os.X_OK): raise ValueError(f'Missing executable server: {executable}') + provenance = weight_selection.prepare_checkpoint(args.weights, args.checkpoint) backbone = sum(p.stat().st_size for p in checkpoint.glob('*.safetensors')) if backbone <= 0: raise ValueError('No backbone safetensors found') @@ -240,10 +248,11 @@ def request_stop(signum: int, _frame: object) -> None: for name in ('hf-cache', 'runtime-cache'): (root / name).mkdir(exist_ok=True) + serve_args = weight_selection.serve_arguments(args.weights, args.serve_args) server_command = [str(executable), 'serve', str(checkpoint), '--host', args.host, '--port', str(args.port), '--name', 'fishaudio/drama-3', '--context', '8192', '--parallel', '1', '--no-drafts', - '--prompt-cache-gib', '0', '--keep-warm', '0', *args.serve_args] + '--prompt-cache-gib', '0', '--keep-warm', '0', *serve_args] environment = { 'HOME': str(Path.home()), 'PATH': f'{root / "venv/bin"}:/usr/bin:/bin:/usr/sbin:/sbin', @@ -254,7 +263,7 @@ def request_stop(signum: int, _frame: object) -> None: 'XDG_CACHE_HOME': str(root / 'runtime-cache'), 'TENSORFOLD_MEMORY_RESERVE_GIB': f'{args.reserve_gib:g}', } - initial = {'backend': args.backend, 'weights_bytes': weights, + initial = {**provenance, 'backend': args.backend, 'weights_bytes': weights, 'initial_memavailable_bytes': available, 'required_memavailable_bytes': int(floor), 'headroom_floor_bytes': HEADROOM, 'abort_threshold_bytes': ABORT_AVAILABLE, diff --git a/recipe/test_weight_selection.py b/recipe/test_weight_selection.py new file mode 100644 index 0000000..d5220da --- /dev/null +++ b/recipe/test_weight_selection.py @@ -0,0 +1,170 @@ +"""Deterministic checkpoint/launcher tests: no network, model, GPU or child process.""" +from __future__ import annotations + +from contextlib import ExitStack +import hashlib +import json +from pathlib import Path +import sys +import tempfile +import unittest +from unittest.mock import patch + +import pilot +import pilot_macos +import weight_selection as selection + +GIB = 1024**3 + + +class WeightSelectionTests(unittest.TestCase): + def setUp(self): + self.stack = ExitStack() + self.addCleanup(self.stack.close) + self.root = Path(self.stack.enter_context(tempfile.TemporaryDirectory())).resolve() + self.recipe = self.root / 'recipe' + (self.recipe / 'checksums').mkdir(parents=True) + self.sources = json.loads((selection.RECIPE / 'sources.json').read_text()) + (self.recipe / 'sources.json').write_text(json.dumps(self.sources)) + self.stack.enter_context(patch.object(selection, 'REPO', self.root)) + self.stack.enter_context(patch.object(selection, 'RECIPE', self.recipe)) + # Real files and real SHA-256, distinct per selection so using the wrong + # manifest cannot accidentally pass. No production weights are opened. + for mode, spec in selection.WEIGHTS.items(): + target = self.root / spec['dir'] + target.mkdir() + entries = [] + for name in ('model.safetensors', 'codec.pth', 'config.json'): + content = f'{mode}:{name}'.encode() + (target / name).write_bytes(content) + entries.append(f'{hashlib.sha256(content).hexdigest()} {name}\n') + (self.recipe / 'checksums' / spec['pins']).write_text(''.join(entries)) + executable = selection.native_path(mode, None) + executable.parent.mkdir(parents=True) + executable.write_bytes(b'not an executable model') + executable.chmod(0o700) + self.download = self.stack.enter_context( + patch.object(selection.subprocess, 'run', side_effect=AssertionError('unexpected download'))) + + + def test_guidance_override_replaces_all_defaults(self): + self.assertEqual(selection.serve_arguments('nvfp4', ['--max-frames', '8']), + ['--cfg-gamma', '1.5', '--cfg-batched', '--max-frames', '8']) + for flag in ('--cfg-gamma', '--cfg-gamma=2', '--cfg-batched', + '--cfg-allow-request-gamma', '--cfg-allow-request-gamma=true'): + arguments = [flag, '2'] if flag == '--cfg-gamma' else [flag] + self.assertEqual(selection.serve_arguments('nvfp4', arguments), arguments) + for mode in ('full', None): + arguments = ['--max-frames', '8'] + self.assertEqual(selection.serve_arguments(mode, arguments), arguments) + self.assertEqual(arguments, ['--max-frames', '8']) + + def test_explicit_paths_override_defaults_but_not_manifest(self): + override = self.root / 'custom-checkpoint' + selection.checkpoint_path('nvfp4', None).rename(override) + binary = self.root / 'custom-native' + self.assertEqual(selection.native_path('nvfp4', binary), binary) + provenance = selection.prepare_checkpoint('nvfp4', override) + self.assertEqual(provenance['checkpoint_path'], str(override)) + self.assertTrue(provenance['checkpoint_verified']) + with self.assertRaises(ValueError): + selection.prepare_checkpoint('full', override) + self.download.assert_not_called() + + def test_missing_override_never_downloads(self): + with self.assertRaises(ValueError): + selection.prepare_checkpoint('nvfp4', self.root / 'missing') + self.download.assert_not_called() + self.assertFalse((self.root / 'missing').exists()) + + def test_existing_missing_or_corrupt_input_is_not_repaired(self): + for mode in selection.WEIGHTS: + target = selection.checkpoint_path(mode, None) + (target / 'model.safetensors').unlink() + with self.subTest(mode=mode, damage='missing'): + with self.assertRaises(ValueError): + selection.prepare_checkpoint(mode, None) + (target / 'model.safetensors').write_bytes(b'corrupt') + with self.subTest(mode=mode, damage='corrupt'): + with self.assertRaises(ValueError): + selection.prepare_checkpoint(mode, None) + self.assertEqual((target / 'model.safetensors').read_bytes(), b'corrupt') + self.download.assert_not_called() + + def test_downloaded_content_is_verified_before_acceptance(self): + target = selection.checkpoint_path('nvfp4', None) + saved = self.root / 'download-fixture' + target.rename(saved) + + def download(_command, **_kwargs): + saved.rename(target) + (target / 'config.json').write_bytes(b'bad download') + + self.download.side_effect = download + with patch.object(selection.shutil, 'which', return_value='/fixture/hf'): + with self.assertRaises(ValueError): + selection.prepare_checkpoint('nvfp4', None) + self.download.assert_called_once() + + def test_legacy_has_no_claim_of_verified_pinned_source(self): + target = self.root / 'operator-checkpoint' + provenance = selection.prepare_checkpoint(None, target) + self.assertEqual(provenance, {'weights_mode': 'explicit', 'checkpoint_path': str(target), + 'checkpoint_source': None, 'checkpoint_manifest': None, + 'checkpoint_verified': False}) + with self.assertRaises(ValueError): + selection.checkpoint_path(None, None) + with self.assertRaises(ValueError): + selection.native_path(None, None) + self.download.assert_not_called() + + def launcher_arguments(self, mode, override=None, binary=None): + lock = self.root / 'gpu.lock' + lock.touch() + args = ['pilot', '--runtime-dir', str(self.root / 'runtime'), + '--evidence-dir', str(self.root / 'evidence'), '--gpu-lock', str(lock)] + if mode is not None: + args += ['--weights', mode] + if override is not None: + args += ['--checkpoint', str(override)] + if binary is not None: + args += ['--native-bin', str(binary)] + return args + + def test_both_launchers_refuse_wrong_selected_override_before_spawn(self): + for launcher in (pilot, pilot_macos): + with self.subTest(launcher=launcher.__name__): + arguments = self.launcher_arguments('nvfp4', selection.checkpoint_path('full', None)) + with patch.object(sys, 'argv', arguments), \ + patch.object(launcher, 'memory_available', return_value=100 * GIB), \ + patch.object(launcher.subprocess, 'Popen') as spawn: + if launcher is pilot: + with self.assertRaises(SystemExit): + launcher.main() + else: + self.assertEqual(launcher.main(), 1) + completion = json.loads((self.root / 'evidence/completion.json').read_text()) + self.assertTrue(completion['group_settled']) + spawn.assert_not_called() + self.download.assert_not_called() + + def test_mac_missing_override_fails_before_spawn_and_invalidates_reused_evidence(self): + arguments = self.launcher_arguments('nvfp4', self.root / 'missing') + evidence = self.root / 'evidence' + evidence.mkdir() + (evidence / 'ready.json').write_text('old readiness') + arguments.append('--evidence-reuse') + with patch.object(sys, 'argv', arguments), \ + patch.object(pilot_macos, 'memory_available', return_value=100 * GIB), \ + patch.object(pilot_macos.subprocess, 'Popen') as spawn: + self.assertEqual(pilot_macos.main(), 1) + spawn.assert_not_called() + self.assertFalse((evidence / 'ready.json').exists()) + completion = json.loads((evidence / 'completion.json').read_text()) + self.assertIsNone(completion['server_pid']) + self.download.assert_not_called() + + + +if __name__ == '__main__': + unittest.main() diff --git a/recipe/weight_selection.py b/recipe/weight_selection.py new file mode 100644 index 0000000..4628fa3 --- /dev/null +++ b/recipe/weight_selection.py @@ -0,0 +1,84 @@ +"""Pinned checkpoint selection shared by the Linux and macOS launchers.""" +from __future__ import annotations + +import hashlib +import json +from pathlib import Path +import shutil +import subprocess +import time + +RECIPE = Path(__file__).resolve().parent +REPO = RECIPE.parent +WEIGHTS = { + 'full': {'source': 'checkpoint', 'dir': 'drama-3', 'pins': 'checkpoint.sha256', 'build': 'tensorfold', + 'serve': ()}, + 'nvfp4': {'source': 'nvfp4_checkpoint', 'dir': 'drama-3-NVFP4-Pruned', 'pins': 'nvfp4.sha256', + 'build': 'tensorfold-nvfp4', 'serve': ('--cfg-gamma', '1.5', '--cfg-batched')}, +} +GUIDANCE_FLAGS = ('--cfg-gamma', '--cfg-allow-request-gamma', '--cfg-batched') + + +def checkpoint_path(choice: str | None, override: Path | None) -> Path: + if override is not None: + return override.resolve() + if choice is None: + raise ValueError('Pass --weights nvfp4 or --weights full, or --checkpoint DIR') + return (REPO / WEIGHTS[choice]['dir']).resolve() + + +def native_path(choice: str | None, override: Path | None) -> Path: + if override is not None: + return override.resolve() + if choice is None: + raise ValueError('The native backend requires --native-bin, or --weights') + return (REPO / 'build' / WEIGHTS[choice]['build'] / 'zig-out/native/bin/tensorfold-native').resolve() + + +def serve_arguments(choice: str | None, arguments: list[str]) -> list[str]: + selected = list(arguments) + if choice and not any(a.split('=', 1)[0] in GUIDANCE_FLAGS for a in selected): + selected[:0] = WEIGHTS[choice]['serve'] + return selected + + +def sha256(path: Path) -> str: + digest = hashlib.sha256() + with path.open('rb') as file: + for block in iter(lambda: file.read(1 << 24), b''): + digest.update(block) + return digest.hexdigest() + + +def prepare_checkpoint(choice: str | None, override: Path | None) -> dict: + """Verify selected pins, downloading only an absent default (never an override). + + Explicit legacy paths have operator-owned provenance, not an implied pinned source. + """ + target = checkpoint_path(choice, override) + provenance = {'weights_mode': choice or 'explicit', 'checkpoint_path': str(target), + 'checkpoint_source': None, 'checkpoint_manifest': None, + 'checkpoint_verified': False} + if choice is None: + return provenance + spec = WEIGHTS[choice] + source = json.loads((RECIPE / 'sources.json').read_text())[spec['source']] + pins = RECIPE / 'checksums' / spec['pins'] + entries = [line.split(None, 1) for line in pins.read_text().splitlines() if line.strip()] + if override is None and not target.exists() and not target.is_symlink(): + hf = shutil.which('hf') + if hf is None: + raise ValueError('the hf command is missing; install it with: python3 -m pip install -U huggingface_hub') + print(f'Downloading {source["repository"]} at {source["revision"][:8]} into {target}', flush=True) + subprocess.run([hf, 'download', source['repository'], '--revision', source['revision'], + '--local-dir', str(target)], check=True) + print(f'Checking {target} against {pins.relative_to(REPO)}', flush=True) + started = time.monotonic() + bad = [name for digest, name in entries if not (target / name).is_file() or sha256(target / name) != digest] + if bad: + raise ValueError(f'{target} does not match {pins}: {", ".join(bad)} missing or changed; no model was ' + f'started. Repair it with: hf download {source["repository"]} --revision ' + f'{source["revision"]} --local-dir {target}, or delete it and run again') + print(f'{target} matches {pins.relative_to(REPO)} ({time.monotonic() - started:.0f} s)', flush=True) + provenance.update(checkpoint_source=source, checkpoint_manifest=str(pins), checkpoint_verified=True) + return provenance diff --git a/results/README.md b/results/README.md index 44786a0..5b5e3f2 100644 --- a/results/README.md +++ b/results/README.md @@ -9,6 +9,7 @@ published; the JSON files retain the measured results and synthetic-input hashes | File | What it holds | |---|---| +| [`macos-nvfp4.json`](macos-nvfp4.json) | Native Metal NVFP4/CFG support on M5 Ultra: 28 HTTP checks, paired/separate PCM identity, registered operator gates with the initial fixture failure retained; no speed or quality-equivalence claim | | [`macos-pr17-repair.json`](macos-pr17-repair.json) | Current full-model Metal cache/lifecycle repair: 48/48 paired corpus captures identical, transition/cancellation/multilane parity, bounded voice-switch RSS; no speed claim | | [`before-after-02/`](before-after-02/README.md) | Reference against native: latency, RTF, first audio, cold start, frame-cap replies | | [`speed-stages.json`](speed-stages.json) | Warm RTF and load time through the five speed stages, with the audio-identity gate | diff --git a/results/macos-nvfp4.json b/results/macos-nvfp4.json new file mode 100644 index 0000000..fa34e90 --- /dev/null +++ b/results/macos-nvfp4.json @@ -0,0 +1,712 @@ +{ + "scope": "Native macOS NVFP4 correctness/support qualification, not a performance benchmark, quality comparison or production deployment", + "platform": { + "os": "macOS Apple silicon", + "observed_gpu": "Apple M5 Ultra", + "observed_gpu_family": "applegpu_g17", + "ram_gib": 256, + "ram_basis": "retained operator inventory; not a capacity qualification" + }, + "compiler": { + "version": "Zig 0.17.0", + "cpu_target": "apple_m1", + "jobs": 2, + "archive_sha256": "b607e9b9234790a008116ae5bdb71c6243b84b9fb42a53a9e70fde41c06c536a" + }, + "checkpoint": { + "repository": "plotarmordev/drama-3-NVFP4-Pruned", + "revision": "494db08481db952f834b2241bce63fb3dfa72554", + "verified_files": 15, + "conversion_or_modification": false + }, + "engine": { + "version": "tensorfold-native 1.1.4", + "base": "3d219d2156d4c25dcfb74370c8703cd89e58aa49", + "upstream_pr": "https://github.com/ashhart/TensorFold/pull/1268", + "native_build_head": "283f624643334b6d1670db01cf4ef61992aff250", + "head_after_test_only_correction": "986477cc879f050f1555cc4e1c5b1d25f06c78de", + "head_note": "The two build heads above are historical qualification heads. Publication-only mail metadata preserves their source; a final CLI-table grouping patch changes capability-list order only and was built/help-checked without replaying GPU qualification.", + "qualified_source_head": "c59b0954129c6ee4e9c9eaa9ebc2fdea3c542f37", + "upstream_source_head": "6585530054e273aa1a1554549585489f2497cb41", + "published_patched_head": "e5e86adb16498b3c172db3a45fafd273c0679fe2", + "published_source_tree": "b143c8a51fb9e53d5e0db384207e32ee877b8594", + "publication_tree_matches_final_source": true, + "native_sha256": "dfb2577c2ad3c543334e9fe8bb3ac7080dd0f7db282455213e8e4dddc0d928be", + "publication_native_sha256": "cd8b3162567dfd98e4305593e8677ee40a340a4ffdb885cee3334db302ef7d1f", + "binary_note": "native_sha256 identifies the live-qualified binary; publication_native_sha256 identifies the later CLI-grouping build, exercised for version/help only.", + "generator_sha256": "1f5167026a043712e69139077ac014bff237aed8240a7bab5f2256797d167140", + "voice_writer_sha256": "d6a2be4eed888c1b77cb3dafeed65947a52ed3ffabbae4bf643e20c4c4f0bb9b", + "patches": [ + { + "file": "0001-drama3-metal-nvfp4-cfg.patch", + "sha256": "f282318f0e8dae43b19d1fa93ec05691696a0352eaeac41f9ffc334690d8b388", + "qualified_mail_sha256": "bafcaaf9156e2971e386547f952addf3e6ea8cd1b667ab91a33e05f55b8abcad" + }, + { + "file": "0002-shared-gate-stress-regression.patch", + "sha256": "0a68a3b66da1991f1a04980cb6958e4acc8856a88efdbe1fe10adc6df588cfc8", + "qualified_mail_sha256": "4e4a86b3ccf789a9f634959a348c3622db4bb38ecc29da27f39295a76a36c15d" + }, + { + "file": "0003-shared-cli-extension-point.patch", + "sha256": "794a5b6008487ec705cadb098eb1c15737da743aee0855b610510506a0ce38db", + "scope": "CLI capability-list ordering only; flag values/parsing, kernels, arithmetic and tolerances unchanged" + } + ] + }, + "build_verification": { + "public_build_entrypoint": "scripts/build-engine.sh --weights nvfp4 BUILD_DIR", + "initial_clean_build": "patch checksum verified; clean apply; native, drama3-generate and drama3-voice built successfully", + "test_only_followup": "second patch applied cleanly with the same git-am options; corrected contract runner built; unchanged production native binary retained, not rebuilt/replayed", + "publication_metadata_proof": "First two checksummed mails applied in a fresh checkout with a Git tree equal to the qualified implementation plus test-only correction. The final CLI-only third patch was applied to a separate checkout of that retained two-patch head; its tree equals final upstream source. Original proof owners exited0, reaped/group absent.", + "final_cli_build": "Final source65855300 native build, version1.1.4 and serve help passed. Original owner exit0/reaped/group absent, no signals or settlement errors; no model/GPU check replayed.", + "integration_scope": "All four Notes-stack merge-tree checks are clean after grouping CFG flags beside voice selection. Reported CUDA conflicts are identical to conflicts against the original base; other PR-pair merges are not claimed clean." + }, + "cpu_verification": { + "weight_selection_tests": 8, + "cfg_host_runner_passed": 20, + "metal_format_guidance_runner_passed": 6, + "independent_nvfp4_reference_tests": 2, + "new_nonlinear_stress_regression": 1, + "kernel_contract_host_tests": 84, + "metal_host_compile_and_tests": "3 passed,2 skipped", + "cuda_host_compile": "passed; no CUDA device execution", + "server_initial": "737 passed,3 skipped,1 incidental busy-reply-count assertion failure; JSON tests5passed", + "server_focused_followup": "16 runner cases passed after removing only incidental count assertion, retaining timeout/redraw assertions; broad suite not replayed", + "lean": "zero new violations on initial feature tree; two inherited violations retained", + "private_verifier_smoke": "original-child natural/graceful settlement and HTTP framing/negative/absolute-deadline checks passed; synthetic responses are not model proof" + }, + "metal_contract": { + "initial_summary": "SUMMARY 14 kernels: 13 pass gates 1-4, 1 fail, 0 known findings; 0 controls: 0 red as planted, 0 missed; 0 limit problems; gates 2138 ms, limits 6587 ms", + "initial_failure": "shared-gate stress saturated the sigmoid, so FP16 accumulation could match the BF16 output; the precision-discrimination gate correctly refused it", + "correction": "test inputs centered the exact FP32 dot at1; independent CPU test proves both reduced accumulators differ after sigmoid/BF16 in all three orders. No production kernel or tolerance change.", + "focused_summary": "SUMMARY 1 kernels: 1 pass gates 1-4, 0 fail, 0 known findings; 0 controls: 0 red as planted, 0 missed; 0 limit problems; gates 159 ms, limits 0 ms", + "coverage": "13 initial successful variants and all runtime resource limits retained; only corrected failed variant rerun. This is not a second full14-variant pass or full multi-family suite." + }, + "live": { + "settings": { + "context": 8192, + "parallel": 1, + "drafts": false, + "prompt_cache_gib": 0, + "keep_warm": 0, + "sampling": "fixed model speech policy; no sampling overrides", + "reference_provenance": "retained attested synthetic catalogs; no natural-person voice" + }, + "counts": { + "total": 28, + "complete_pcm": 15, + "early_stream_cancellation": 1, + "expected_http_400": 12, + "failed_checks": 0 + }, + "phases": [ + { + "phase": "paired", + "passed": true, + "request_count": 8, + "checks": { + "a_repeat_exact": true, + "a_after_b_exact": true, + "a_stream_whole_exact": true, + "early_cancellation": true, + "a_after_cancellation_exact": true, + "request_count_exact": true + }, + "hashes": { + "default": "04d31906e89b75d9545ce0358e16c734e4b1e86ce2e389c6da0a0f32768b651f", + "synthetic-a": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "synthetic-b": "3e0a49f287e6796570862dd1f810d82013d22409d5ce167cf525ca682b272c4a" + }, + "requests": [ + { + "label": "default", + "status": 200, + "passed": true, + "bytes": 331776, + "sha256": "04d31906e89b75d9545ce0358e16c734e4b1e86ce2e389c6da0a0f32768b651f", + "frames": 81.0, + "audio_duration_s": 3.7616326530612243, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-repeat", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "b", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "3e0a49f287e6796570862dd1f810d82013d22409d5ce167cf525ca682b272c4a", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-after-b", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-stream", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-sample-rate": "44100", + "x-channels": "1", + "x-bit-depth": "16", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-cancel", + "status": 200, + "passed": true, + "bytes": 4096, + "sha256": "fb1a96188111e39464eba33fbf96426bcc72f919d6756f1481c82dbacb2bb890", + "frames": 1.0, + "audio_duration_s": 0.046439909297052155, + "finish_reason": null, + "complete": false, + "cancelled": true, + "response_headers": { + "content-type": "audio/pcm", + "x-sample-rate": "44100", + "x-channels": "1", + "x-bit-depth": "16", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-after-cancel", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + } + ], + "settlement": { + "exit_code": 0, + "reaped": true, + "settled": true, + "escalated": false, + "settlement_errors": [] + }, + "stop": "planned SIGTERM to original native child" + }, + { + "phase": "separate", + "passed": true, + "request_count": 4, + "checks": { + "paired_hashes_available": true, + "paired_separate_exact_default": true, + "paired_separate_exact_synthetic-a": true, + "paired_separate_exact_synthetic-b": true, + "a_stream_whole_exact": true, + "request_count_exact": true + }, + "hashes": { + "default": "04d31906e89b75d9545ce0358e16c734e4b1e86ce2e389c6da0a0f32768b651f", + "synthetic-a": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "synthetic-b": "3e0a49f287e6796570862dd1f810d82013d22409d5ce167cf525ca682b272c4a" + }, + "requests": [ + { + "label": "default", + "status": 200, + "passed": true, + "bytes": 331776, + "sha256": "04d31906e89b75d9545ce0358e16c734e4b1e86ce2e389c6da0a0f32768b651f", + "frames": 81.0, + "audio_duration_s": 3.7616326530612243, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "b", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "3e0a49f287e6796570862dd1f810d82013d22409d5ce167cf525ca682b272c4a", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1.5" + } + }, + { + "label": "a-stream", + "status": 200, + "passed": true, + "bytes": 327680, + "sha256": "c8769b5ffac317dde8c3a338ce81a6311fa3a6588779b865f5873b661105b837", + "frames": 80.0, + "audio_duration_s": 3.7151927437641725, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-sample-rate": "44100", + "x-channels": "1", + "x-bit-depth": "16", + "x-cfg-gamma": "1.5" + } + } + ], + "settlement": { + "exit_code": 0, + "reaped": true, + "settled": true, + "escalated": false, + "settlement_errors": [] + }, + "stop": "planned SIGTERM to original native child" + }, + { + "phase": "request", + "passed": true, + "request_count": 16, + "checks": { + "gamma_1_repeat_exact": true, + "gamma_2_repeat_exact": true, + "request_count_exact": true + }, + "hashes": { + "synthetic-a": { + "gamma-1": "a44e67300b185e27e96fc8a2e668c7d33d96b24dd394c56e2b202fbae3e8ea72", + "gamma-2": "78453aaf336934f307d0306f1f253cc6f342973d223a3cc647e15230d6f155d4" + } + }, + "requests": [ + { + "label": "gamma-1", + "status": 200, + "passed": true, + "bytes": 307200, + "sha256": "a44e67300b185e27e96fc8a2e668c7d33d96b24dd394c56e2b202fbae3e8ea72", + "frames": 75.0, + "audio_duration_s": 3.4829931972789114, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1" + } + }, + { + "label": "gamma-1-repeat", + "status": 200, + "passed": true, + "bytes": 307200, + "sha256": "a44e67300b185e27e96fc8a2e668c7d33d96b24dd394c56e2b202fbae3e8ea72", + "frames": 75.0, + "audio_duration_s": 3.4829931972789114, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "1" + } + }, + { + "label": "gamma-2", + "status": 200, + "passed": true, + "bytes": 315392, + "sha256": "78453aaf336934f307d0306f1f253cc6f342973d223a3cc647e15230d6f155d4", + "frames": 77.0, + "audio_duration_s": 3.575873015873016, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "2" + } + }, + { + "label": "gamma-2-repeat", + "status": 200, + "passed": true, + "bytes": 315392, + "sha256": "78453aaf336934f307d0306f1f253cc6f342973d223a3cc647e15230d6f155d4", + "frames": 77.0, + "audio_duration_s": 3.575873015873016, + "finish_reason": "stop", + "complete": true, + "cancelled": false, + "response_headers": { + "content-type": "audio/pcm", + "x-finish-reason": "stop", + "x-cfg-gamma": "2" + } + }, + { + "label": "bad-gamma-0", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "X-Cfg-Gamma must be a number from 1 (unguided) to 5", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "bad-gamma-1", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "X-Cfg-Gamma must be a number from 1 (unguided) to 5", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "bad-gamma-2", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "X-Cfg-Gamma must be a number from 1 (unguided) to 5", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "bad-gamma-3", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "X-Cfg-Gamma must be a number from 1 (unguided) to 5", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "bad-gamma-4", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "X-Cfg-Gamma must be a number from 1 (unguided) to 5", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-temperature", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "temperature is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-top_p", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "top_p is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-top_k", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "top_k is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-repetition_penalty", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "repetition_penalty is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-seed", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "seed is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-stage_params", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "stage_params is not supported by this model", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + }, + { + "label": "unsupported-min_p", + "status": 400, + "passed": true, + "finish_reason": null, + "complete": true, + "cancelled": false, + "pcm_bytes": 0, + "error_response": { + "error": { + "message": "unknown speech request field: min_p", + "type": "invalid_request_error" + } + }, + "response_headers": { + "content-type": "application/json" + } + } + ], + "settlement": { + "exit_code": 0, + "reaped": true, + "settled": true, + "escalated": false, + "settlement_errors": [] + }, + "stop": "planned SIGTERM to original native child" + } + ] + }, + "settlement": { + "operator_owners": "both original children reaped and groups absent, no signals or settlement errors", + "live_owner": "exit0/reaped/group absent, no signals or settlement errors", + "native_servers": 3, + "native_exit_codes": [ + 0, + 0, + 0 + ], + "planned_stops_only": true, + "guard_failures": 0, + "cooperative_holds_released": true, + "production_changes": false + }, + "limitations": [ + "Tested on M5 Ultra only; M1-M4 runtime and smaller-memory Macs unqualified.", + "No CUDA execution or cross-backend PCM comparison.", + "No WER, speaker-quality, perceptual-equivalence or speed claim.", + "No full multi-family/rotary Metal suite or offline M1/M2 compiler receipts.", + "Public Mac launcher selection/verification has CPU coverage; live engine ran directly under the existing cooperative guard to avoid nested flock ownership.", + "Early close and exact follow-up PCM are observed; engine-side cancellation was not independently instrumented.", + "HTTP400 precedes PCM but does not independently prove absence of engine submission.", + "No whole-host GPU-idleness/exclusivity claim; no peak process/device-memory capacity measurement.", + "Full-model and CUDA pins remain unchanged; completed earlier experiments were not rerun." + ], + "retained_preparation_failures": [ + "Initial Zig integration/contract compilation errors were corrected; only failed or newly changed targets were rerun.", + "Initial broad server test retained the incidental scheduler-count failure; focused consumer-visible deadline checks passed after removal.", + "Initial underpowered GPU stress fixture remains recorded above; it was not relabeled as a passing run." + ] +} diff --git a/scripts/build-engine.sh b/scripts/build-engine.sh index ea63b29..2565afa 100755 --- a/scripts/build-engine.sh +++ b/scripts/build-engine.sh @@ -5,8 +5,9 @@ # # --weights full (the default) upstream TensorFold at engine/UPSTREAM with the opt-in patches in # engine/patches: the native server and drama3-generate, in build/tensorfold -# --weights nvfp4 the engine at engine/nvfp4/UPSTREAM, which serves the NVFP4 pruned weights -# (plotarmordev/drama-3-NVFP4-Pruned), unpatched: the native server, in build/tensorfold-nvfp4 +# --weights nvfp4 CUDA: engine/nvfp4/UPSTREAM, unpatched native server. +# Metal: engine/nvfp4/macos/UPSTREAM plus its patches, native server and voice tools. +# Both use plotarmordev/drama-3-NVFP4-Pruned, in build/tensorfold-nvfp4. # # Environment: # ZIG Zig 0.17.0 executable (default: zig on PATH) @@ -14,13 +15,14 @@ # JOBS parallel build jobs (default: Zig's choice) # # Every run checks out the pinned commit and re-applies the whole series, so the result does not depend on what an -# earlier run left behind. Output: BUILD_DIR/zig-out/native/bin/tensorfold-native, and with the full weights also -# BUILD_DIR/zig-out/bin/drama3-generate. +# earlier run left behind. Output: BUILD_DIR/zig-out/native/bin/tensorfold-native, plus the selected voice tools +# under BUILD_DIR/zig-out/bin. set -euo pipefail repo=$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd) zig=${ZIG:-zig} nvcc=${NVCC:-nvcc} +platform=$(uname -s) die() { echo "build-engine: $*" >&2 @@ -37,9 +39,22 @@ elif [[ ${1:-} == --weights=* ]]; then shift fi (($# <= 1)) || die "usage: scripts/build-engine.sh [--weights full|nvfp4] [BUILD_DIR]" +patch_dir= case $weights in - full) pin=$repo/engine/UPSTREAM build=${1:-$repo/build/tensorfold} targets=(native drama3-generate) ;; - nvfp4) pin=$repo/engine/nvfp4/UPSTREAM build=${1:-$repo/build/tensorfold-nvfp4} targets=(native) ;; + full) + pin=$repo/engine/UPSTREAM build=${1:-$repo/build/tensorfold} + patch_dir=$repo/engine/patches + targets=(native drama3-generate) + ;; + nvfp4) + pin=$repo/engine/nvfp4/UPSTREAM build=${1:-$repo/build/tensorfold-nvfp4} + targets=(native) + if [[ $platform == Darwin ]]; then + pin=$repo/engine/nvfp4/macos/UPSTREAM + patch_dir=$repo/engine/nvfp4/macos/patches + targets+=(drama3-generate drama3-voice) + fi + ;; *) die "--weights must be full or nvfp4, not '$weights'" ;; esac @@ -49,19 +64,27 @@ esac shopt -s nullglob patches=() -if [[ $weights == full ]]; then - patches=("$repo"/engine/patches/*.patch) - ((${#patches[@]})) || die "engine/patches has no *.patch files" +if [[ -n $patch_dir ]]; then + patches=("$patch_dir"/*.patch) + ((${#patches[@]})) || die "$patch_dir has no *.patch files" + if command -v sha256sum >/dev/null; then + (cd "$patch_dir" && sha256sum -c SHA256SUMS) || die "patch checksum verification failed" + else + (cd "$patch_dir" && shasum -a 256 -c SHA256SUMS) || die "patch checksum verification failed" + fi fi command -v git >/dev/null || die "git not found" zig=$(command -v "$zig") || die "Zig not found; set ZIG" -# macOS builds the Metal engine; the CUDA option is not declared there, so do not pass it. -if [[ $(uname -s) == Darwin ]]; then - nvcc_args=() +# macOS builds the Metal engine without requiring a CUDA compiler. +if [[ $platform == Darwin ]]; then + build_args=() + if [[ $weights == nvfp4 ]]; then + build_args=(-Dcpu=apple_m1) + fi else nvcc=$(command -v "$nvcc") || die "nvcc not found; set NVCC" - nvcc_args=(-Dnvcc="$nvcc") + build_args=(-Dnvcc="$nvcc") fi if [[ -e $build ]]; then @@ -88,13 +111,13 @@ have=$("$zig" version) [[ $have == "$want" ]] || die "Zig $have found; the engine pins $want" cd "$build" -"$zig" build "${targets[@]}" ${nvcc_args[@]+"${nvcc_args[@]}"} --summary all ${JOBS:+-j"$JOBS"} +"$zig" build "${targets[@]}" ${build_args[@]+"${build_args[@]}"} --summary all ${JOBS:+-j"$JOBS"} echo "engine ($weights weights): $TENSORFOLD_COMMIT + ${#patches[@]} patches = $(git rev-parse HEAD)" artifacts=(zig-out/native/bin/tensorfold-native) -if [[ $weights == full ]]; then - artifacts+=(zig-out/bin/drama3-generate) -fi +for target in "${targets[@]:1}"; do + artifacts+=("zig-out/bin/$target") +done if command -v sha256sum >/dev/null; then sha256sum "${artifacts[@]}" else