drama: add native Metal NVFP4 and reference CFG support

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.

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`. |

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);

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.

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,

zig/build/kernel_contract.zig

@@ -36,6 +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_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.

zig/kernels/metal/drama3/guidance.metal

@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+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]);
+}

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);
+  }
+}

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");
 };

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");

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 },

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,