Nemotron 3.5 Transcribe: audio input, the 65,536-token window, an hour of audio, MTP drafts

docs/api.md

@@ -129,6 +129,8 @@ server reads images for GLM-5.3-Flash, Qwen3.8-27B and Gemma 4 31B on Metal, wit
 user messages ([video](vision.md#video-on-the-native-server)).
 See [image input](vision.md) for data URLs, public image URLs, limits and cache behavior.
 The native binary refuses `--vision` for other models and answers their image parts with HTTP 400.
+The native CUDA server reads audio for Nemotron 3.5 Transcribe: `input_audio` and `audio_url` (data URL) parts with
+16 kHz mono WAV in user messages ([recipe](recipes/nemotron-3.5-transcribe.md#audio-in-a-request)).
 Unsupported image input, audio, video and non-text output requests receive HTTP 400.
 
 ## Decisions

docs/kernel-contract.md

@@ -183,6 +183,11 @@ CUDA also checks Dory's activation rows (`act_l2.cuh`): the scaled sigmoid outpu
 each row then L2-normed, against float64's ranges at their bf16 points and expf's ulps (four controls, one point each).
 CUDA also checks `dense_rows`' scaled epilogue (Dory's logits, fp32(bf16(x . w)) times a per-column scale), bit-equal to its
 unscaled tile's outputs rounded and scaled on the host.
+CUDA also checks Nemotron 3.5 Transcribe's audio encoders (`nemo_audio.cu`): the ASR encoder's q/k LayerNorm with
+RoPE on 64 of 128 channels and the Sortformer's RoPE alone (each output within float64's bound at its three bf16
+points), attention over a window's valid keys (core/attention_range.zig's range, P as bf16), and the Sortformer head's
+im2col, padding and relu rows, exact (five controls: a one-pass variance, a key short of the window, a padding row read
+or kept, and a relu by `fmaxf`).
 HIP: the RMS norm rows, a wave a row (fp32, fp16 and bf16 out) and a block a row (fp16 and bf16 out).
 
 RADIO's fused bias, activation and residual (`tf_vit_bias_act`) replaces `tf_vit_bias` then
@@ -203,6 +208,9 @@ CUDA also checks the affine-3 experts' split decode pair, bit-equal to the pair
 one dropping a last round of K groups short of the warps, one chaining the groups in reverse).
 CUDA also checks Flash Next's NVFP4 experts on prompt tiles with the decode form's sums (SwiGLU gate-up and fp32
 down, each bit-equal to the decode instance it replaces; the scaled-code prompt tile is the planted control).
+CUDA also checks the bf16 experts (`experts_bf16.cu`, a checkpoint's own bf16 weights: Nemotron 3.5 Transcribe's routed
+and shared experts): relu^2 up, fp32 down and bf16 down, in the decode and prompt forms, against float64 (two controls:
+the fp32 down launched one 32-input group short, and a down summing in bf16).
 CUDA also checks Nemotron's pocket (`nemotron_ops.cu`, the learned neurons the living-weights learner adds beside the
 shared expert): relu^2 up into act rows and down added into the shared expert's slot, a block row a row and tiled, the
 down to fp32 and bf16, against float64, the neurons past the count NaN so a read of one shows (two controls: the up's

docs/recipes/README.md

@@ -9,6 +9,7 @@ For agent-assisted contributions, read the [TensorFold contributor skill](SKILL.
 | --- | --- |
 | Nemotron 3.5 Lightning | [MLX](nemotron-3.5.md), [Intel Arc, experimental](xpu-arc.md) |
 | Nemotron 3.5 Super VL | [Native CUDA on one or two GB10s, images](nemotron-3.5-super-vl.md) |
+| Nemotron 3.5 Transcribe | [Native CUDA on one GB10, the bf16 release, audio input](nemotron-3.5-transcribe.md) |
 | Qwen3.8-27B | [MLX, quantization and CUDA](qwen3.8-27b.md) |
 | Qwen3.5-2B | [Native Metal preview, tied affine 4-bit](qwen3.5-2b-native.md) |
 | Qwen3.8 Flash Next | [MLX prefill and CUDA](qwen3.8-flash-next.md), [native CUDA](flash-next-cuda.md), [CUDA images](flash-next-vision.md), [exact SSD residency on Metal](fluid-subgraph-residency.md) |

docs/recipes/nemotron-3.5-transcribe.md

@@ -0,0 +1,107 @@
+# Nemotron 3.5 Transcribe
+
+NVIDIA's Nemotron 3.5 Transcribe 30B-A4B (`nemo_speechlm`) runs on the native CUDA engine from NVIDIA's bf16 release,
+on one GB10. It is a speech LLM: an audio encoder turns a recording into rows that stand in the prompt where the
+chat template writes `<|audio|>`, and a Nemotron 3.5 Lightning language model (Mamba-2, attention and MoE blocks,
+128 routed experts with relu^2, a shared expert, one MTP layer) writes the transcript, translation or answer. The
+language model is the Nemotron family (`zig/src/families/nemotron/`, the checkpoint's names in `cuda_salm.zig`); the
+encoder is `zig/src/families/nemo_audio/`, with its kernels in `zig/kernels/cuda/nemo_audio.cu`.
+
+## Checkpoint
+
+`nvidia/NVIDIA-Nemotron-3.5-Transcribe-30B-A4B` at revision `c421d6acdb436edd7ae00565f6414631fb9bd56d`: one
+`model.safetensors` of 63 GiB in NeMo's layout, every tensor bf16 except the Mamba layers' fp32 parameters, the
+routers' fp32 biases and the CTC timestamp head. The engine reads these tensors as stored; nothing is converted
+ahead of time. The routed experts stay in bf16 on a grouped kernel of their own (`zig/kernels/cuda/experts_bf16.cu`).
+
+Not read: NVIDIA's FP8 and NVFP4 variants, the DFlash2 draft model, and the CTC timestamp head.
+
+## Run on one GB10
+
+```bash
+tensorfold serve /models/NVIDIA-Nemotron-3.5-Transcribe-30B-A4B --name transcribe
+```
+
+The weights and the encoder take 64.7 GiB. The default window is the card's trained 65,536 tokens (config.json's
+`llm_config` names the backbone's 262,144, which the speech model was not trained to; `--context` above 65,536 is
+refused). At that window the server admits 8 streams, the engine's own at 0.50 GiB and 7 more taking 512-row pages
+(3.09 GiB in all) as they need them, and keeps a prompt cache of 64 states (3.0 GiB) and 13.0 GiB of copied rows.
+
+The chat template turns thinking on by default. The card asks for transcription, translation and language
+identification with thinking off and temperature 0, and for audio questions with thinking on, temperature 0.6 and
+top_p 0.95. Send `chat_template_kwargs: {"enable_thinking": false}` per request, or start the server with
+`--no-thinking`.
+
+## Audio in a request
+
+`/v1/chat/completions` (and `/v1/completions` given `messages`) takes audio in user messages as OpenAI's
+`input_audio` part or as vLLM's `audio_url` part with a data URL:
+
+```json
+{"messages": [{"role": "user", "content": [
+   {"type": "text", "text": "Transcribe the following audio."},
+   {"type": "input_audio", "input_audio": {"data": "<base64 WAV>", "format": "wav"}}]}],
+ "temperature": 0, "chat_template_kwargs": {"enable_thinking": false}}
+```
+
+- A clip is a WAV file: 16 kHz, mono, 16-bit PCM or 32-bit float, at least 10 ms. Other rates, channels and formats
+  get HTTP 400 naming the clip.
+- A request holds up to 8 clips and up to an hour of audio in all (12.5 rows a second, 45,000 rows), and no more
+  than its window leaves after the text. Chat bodies may be 160 MiB on this model, enough for an hour of 16-bit
+  WAV in base64; other routes keep the 32 MiB limit.
+- `/v1/messages` refuses audio: Anthropic's API has no audio block.
+- Audio parts in assistant, system or tool messages are refused, as is text that spells `<|audio|>` with no clip for it.
+- The card's [prompt guide](https://huggingface.co/nvidia/NVIDIA-Nemotron-3.5-Transcribe-30B-A4B#prompt-guide)
+  lists the instructions for written or verbatim form, context terms, speaker tags (`<spk:0>` ...), translation and
+  language identification; they are plain text beside the clip.
+
+The encoder runs on each request's prompt pass: log-mel features on the host, the 32-layer Transformer ASR encoder
+and the streaming Sortformer speaker encoder on the GPU, in 30-second windows, fused and projected to the language
+model's 2,688 columns. A clip's rows are keyed by a hash of its samples in the prompt cache, so a repeated clip (a
+second question about the same recording) resumes after it instead of encoding it again. `tensorfold audio-check
+MODEL WAV OUT.safetensors` writes every encoder stage's rows for one clip, for comparison with NeMo's.
+
+## Drafts
+
+The checkpoint's MTP layer drafts by default. Its head reads the shifted prompt with the clip's audio rows in place,
+as NeMo trains it. `--mtp-drafts` and `--mtp-confidence` set it as for the other Nemotron checkpoints;
+`"draft": false` decodes serially. Copied continuations stay off for this model. The head adds 2.9 GiB (64.73 GiB against 61.85 with
+`--no-drafts`).
+
+## Exactness and speed
+
+Measured on one GB10 at the default 65,536-token window, greedy, thinking off, "Transcribe the following audio.":
+
+- **Against NeMo's model in bf16:** on five LibriSpeech clips (3.5, 5.9, 34 and 130 seconds, and a 19.5-second
+  clip of two speakers), four replies are the reference's tokens, the 130-second clip's 379 among them. The 34-second
+  clip writes "at 'em" for the reference's "at em", a near tie that the engine's bf16 logits settle the other way.
+- **Encoder stages:** at every stage the engine's rows are as far from NeMo's fp32 rows as NeMo's own bf16 rows are,
+  on eight clips; the speaker features are identical on all of them.
+- **Drafted against `--no-drafts`:** byte-identical replies on 11, 30 and 60 minutes of audio.
+- **Prompt cache:** a server with `--prompt-cache-gib 0` gave the same replies as the default one to every check
+  below: two instructions after the same 11-minute clip, a question after a transcribed clip, seeded sampling,
+  speaker tags, translation to German and French, language identification, audio questions, thinking on, a two-turn
+  chat and the Anthropic route with text. A repeated clip resumed 261 of 267 prompt tokens and answered its first
+  token in 0.067 s instead of 0.381 s.
+- **Concurrency:** four long clips at once (11, 11 and 15 minutes and 130 seconds) took 81.8 s together, each reply
+  identical to its solo run.
+
+Long recordings, LibriSpeech test-clean read end to end without repeats (WER on the normalized text):
+
+| Audio | Prompt tokens | Reply tokens | WER | First token | MTP drafts | `--no-drafts` |
+| --- | --- | --- | --- | --- | --- | --- |
+| 11 min | 8,730 | 2,061 | 1.13% | 13.3 s | 40.4 s, 76.1 tok/s | 73.7 s, 34.1 tok/s |
+| 30 min | 22,374 | 5,238 | 1.27% | 33.9 s | 104.2 s, 74.8 tok/s | 190.8 s, 33.4 tok/s |
+| 60 min | 44,872 | 11,084 | 1.18% | 70.2 s | 228.9 s, 70.0 tok/s | 411.2 s, 32.4 tok/s |
+
+The first token waits for the encoder and the whole prompt pass. The drafts kept 72% to 76% of their proposals
+(9,748 of 13,142 on the hour). The 130-second clip answered its first token in 2.32 s. The encoder's attention is a
+plain per-window kernel, not yet a tiled one. The harness's quality suite was not run on this build.
+
+## Not served yet
+
+- Word timestamps and diarization segments (the card's `ctc_timestamps`, from its CTC head): the head stays
+  unloaded.
+- Audio URLs other than data URLs, and formats other than WAV.
+- The DFlash2 draft model, and the FP8 and NVFP4 variants.
+- More than one GPU: the model fits one GB10, and the two-rank expert split is Super VL's alone.

zig/build/cuda.zig

@@ -32,6 +32,8 @@ fn family(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builtin
     nemotron.addImport("nemotron_draft_ids", draft_ids);
     const radio = radioModule(b, target, optimize, cuda, core, lanes);
     nemotron.addImport("radio", radio);
+    // Nemotron 3.5 Transcribe's audio encoder: the rows its prompt's <|audio|> positions take
+    nemotron.addImport("nemo_audio", b.createModule(.{ .root_source_file = b.path("zig/src/families/nemo_audio/audio.zig"), .target = target, .optimize = optimize, .link_libc = true, .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = core } } }));
     nemotron.addImport("living", living);
     const qwen35 = b.createModule(.{ .root_source_file = b.path("zig/src/families/qwen3_5/cuda.zig"), .target = target, .optimize = optimize, .link_libc = true });
     qwen35.addImport("living", living);

zig/build/cuda_family_host.zig

@@ -27,6 +27,9 @@ pub fn steps(b: *std.Build, cuda: *std.Build.Module, mods: anytype, step: *std.B
     b.step("test-qwen35-formats-host", "MLX, ModelOpt and EXL3 admission and actual embedding dispatch without devices").dependOn(&formats_run.step);
     step.dependOn(&formats_run.step);
     @import("glm_contract_host.zig").host(b, cuda, mods.core, step);
+    const nemo_audio = b.addRunArtifact(b.addTest(.{ .root_module = mods.nemotron.import_table.get("nemo_audio").? }));
+    b.step("test-nemo-audio-host", "Transcribe's audio config, FFT, bf16 rounding and Sortformer FIFO without devices").dependOn(&nemo_audio.step);
+    step.dependOn(&nemo_audio.step);
     @import("dory.zig").host(b, @import("dory.zig").module(b, host, .debug, cuda, mods.core, mods.lanes), step);
     const attention = b.addRunArtifact(b.addTest(.{ .root_module = mods.qwen35, .filters = &.{"Qwen attention path:"} }));
     b.step("test-qwen35-attention-path-host", "Contiguous and paged CUDA attention selection and graph identity without devices").dependOn(&attention.step);

zig/build/cuda_kernels.zig

@@ -60,6 +60,7 @@ pub const kernels = [_]Kernel{
     .{ .name = "exlr8", .flags = &.{ "-O3", "--fmad=false" } }, // ours: EXLR8 lane projections, the Metal kernels' order
     .{ .name = "exlr8_dense", .flags = &.{ "-O3", "--fmad=false" } }, // ours: EXLR8 dense matrices by row, the Metal order
     .{ .name = "nvfp4_experts_tile", .flags = &.{"-O3"}, .min_sm = 80 },
+    .{ .name = "experts_bf16", .flags = &.{"-O3"} }, // ours: grouped bf16 experts on the plan, mma fragment blocks
     .{ .name = "nvfp4_experts", .flags = &.{"-O3"} }, // nvfp4/experts.cu's device code, the NVFP4 expert kernels
     .{ .name = "rank_sum", .flags = &.{"-O3"} }, // ours: the rank-order sum of a reduce-scatter (rank_scatter.zig)
     .{ .name = "lane_gemv", .flags = &.{"-O3"} }, // ours: qmm_group's arithmetic, a column tile's K slices in one CTA
@@ -84,6 +85,7 @@ pub const kernels = [_]Kernel{
     .{ .name = "lfm2", .flags = glue }, // ours: LFM2's short convolution, QK norm and rotary, gates and label rows
     .{ .name = "vit", .flags = glue }, // ours: vision towers' LayerNorm, Linear epilogues, activations, unshuffle
     .{ .name = "moss_image", .flags = glue }, // ours: MOSS-Image's modulated norm, q/k rotary, situ, gated add
+    .{ .name = "nemo_audio", .flags = glue }, // ours: Nemotron 3.5 Transcribe's audio encoders (q/k norm and RoPE, attention, Sortformer glue)
     .{ .name = "tree_round", .flags = &.{"-O3"} }, // ours: verified tree rounds for every family (cuda/tree_round.zig)
     .{ .name = "act_rows", .flags = &.{"-O3"} }, // ours: SwiGLU, output gates, group sums (cuda/act_rows.zig)
     .{ .name = "gdn_prep", .flags = &.{"-O3"} }, // ours: Gated DeltaNet conv, norms, g and beta (cuda/gdn_prep.zig)

zig/kernels/cuda/contract_controls.cu

@@ -352,6 +352,9 @@ TF_CTL_SWIGLU_L2(tf_ctl_swiglu_l2_drop_us, tf_act_l2::GS | tf_act_l2::PROD)
 TF_CTL_SWIGLU_L2(tf_ctl_swiglu_l2_drop_prod, tf_act_l2::GS | tf_act_l2::US)
 #undef TF_CTL_SWIGLU_L2
 
+// nemo_audio.cu's controls: tf_na_qk_rope with a one-pass variance (gate 3), its relu as fmaxf (gate 1).
+#include "contract_nemo_audio_controls.cuh"
+
 // Gate 1: tf_pocket_in (a warp a neuron, act rows 64 wide) with relu as fmaxf(s, 0), which turns a NaN sum into 0.
 extern "C" __global__ void __launch_bounds__(128) tf_ctl_pocket_in_fmax(const __nv_bfloat16* y, int d, const float* up,
                                                                        const int* used, float* act, int rows) {
@@ -538,3 +541,26 @@ TF_CTL_EXL3_GROUPED(tf_ctl_exl3_grouped_old_2, 2, tf_exl3::PLAIN)
 TF_CTL_EXL3_GROUPED(tf_ctl_exl3_grouped_skip, 2, tf_exl3::SKIP)
 TF_CTL_EXL3_GROUPED(tf_ctl_exl3_grouped_stale, 2, tf_exl3::STALE)
 #undef TF_CTL_EXL3_GROUPED
+
+// Gate 4: experts_bf16.cu's fp32 down with each pair's sum rounded to bf16 at every input, a thread a column.
+extern "C" __global__ void tf_ctl_bf16_experts_bf16_sums(const __nv_bfloat16* X, int x_stride, const uint4* W, int KG, int NB,
+                                                         const int* items, const int* counts, const int* members, float* out,
+                                                         int N) {
+    const int col = blockIdx.x * blockDim.x + threadIdx.x;
+    if (col >= N) return;
+    const int cb = col / 32, j = (col % 32) / 8;
+    for (int it = blockIdx.y; it < counts[0]; it += gridDim.y) {
+        const int e = items[3 * it], first = items[3 * it + 1], cnt = items[3 * it + 2];
+        const __nv_bfloat16* blk = reinterpret_cast<const __nv_bfloat16*>(W + (static_cast<int64_t>(e) * NB + cb) * KG * 128);
+        for (int i = 0; i < cnt; ++i) {
+            const int pr = members[first + i];
+            float s = 0.0f;
+            for (int k = 0; k < KG * 32; ++k) {
+                const int q = k % 16, lane = 4 * (col % 8) + q / 4;
+                const float w = __bfloat162float(blk[(static_cast<int64_t>(k / 32) * 128 + 32 * j + lane) * 8 + 4 * (k % 32 / 16) + q % 4]);
+                s = __bfloat162float(__float2bfloat16_rn(s + tf_ctl_x(X, static_cast<int64_t>(pr) * x_stride + k) * w));
+            }
+            out[static_cast<int64_t>(pr) * N + col] = s;
+        }
+    }
+}

zig/kernels/cuda/contract_nemo_audio_controls.cuh

@@ -0,0 +1,56 @@
+// nemo_audio.cu's planted controls (Nemotron 3.5 Transcribe's audio encoder); only the contract controls include this.
+#pragma once
+#include <cuda_bf16.h>
+
+// tf_na_qk_rope with a one-pass variance, E[x^2] - mean^2: rows far from zero lose their digits, so gate 3 goes red.
+__device__ __forceinline__ float tf_ctl_na_bf(__nv_bfloat16 v) { return __bfloat162float(v); }
+__device__ __forceinline__ float tf_ctl_na_rbf(float v) { return __bfloat162float(__float2bfloat16_rn(v)); }
+__device__ __forceinline__ float tf_ctl_na_sum(float v) {
+    for (int o = 16; o > 0; o >>= 1) v = __fadd_rn(v, __shfl_xor_sync(0xffffffffu, v, o));
+    return v;
+}
+
+extern "C" __global__ void __launch_bounds__(32) tf_ctl_na_qk_onepass(__nv_bfloat16* qkv, const __nv_bfloat16* qw, const __nv_bfloat16* qb,
+                                                                          const __nv_bfloat16* kw, const __nv_bfloat16* kb, const __nv_bfloat16* cs,
+                                                                          const __nv_bfloat16* sn, int H, int D, int R, float eps) {
+    __shared__ float x[128];
+    const int which = blockIdx.x & 1;
+    const int h = (blockIdx.x >> 1) % H;
+    const int t = (blockIdx.x >> 1) / H;
+    __nv_bfloat16* row = qkv + ((static_cast<int64_t>(t) * 3 + which) * H + h) * D;
+    const int lane = threadIdx.x;
+    for (int i = lane; i < D; i += 32) x[i] = tf_ctl_na_bf(row[i]);
+    __syncwarp();
+    const __nv_bfloat16* w = which ? kw : qw;
+    const __nv_bfloat16* b = which ? kb : qb;
+    float acc = 0.0f, sq = 0.0f;
+    for (int i = lane; i < D; i += 32) {
+        acc = __fadd_rn(acc, x[i]);
+        sq = __fadd_rn(sq, __fmul_rn(x[i], x[i]));
+    }
+    const float mean = __fdiv_rn(tf_ctl_na_sum(acc), static_cast<float>(D));
+    const float var = __fsub_rn(__fdiv_rn(tf_ctl_na_sum(sq), static_cast<float>(D)), __fmul_rn(mean, mean));
+    const float rstd = __fdiv_rn(1.0f, __fsqrt_rn(__fadd_rn(fmaxf(var, 0.0f), eps)));
+    float y[4];
+    for (int i = lane, k = 0; i < D; i += 32, ++k) y[k] = tf_ctl_na_rbf(__fadd_rn(__fmul_rn(__fmul_rn(__fsub_rn(x[i], mean), rstd), tf_ctl_na_bf(w[i])), tf_ctl_na_bf(b[i])));
+    __syncwarp();
+    for (int i = lane, k = 0; i < D; i += 32, ++k) x[i] = y[k];
+    __syncwarp();
+    const int half = R / 2;
+    for (int i = lane; i < D; i += 32) {
+        float v = x[i];
+        if (i < R) {
+            const float c = tf_ctl_na_bf(cs[static_cast<int64_t>(t) * R + i]);
+            const float s = tf_ctl_na_bf(sn[static_cast<int64_t>(t) * R + i]);
+            const float r = i < half ? -x[i + half] : x[i - half];
+            v = __fadd_rn(tf_ctl_na_rbf(__fmul_rn(v, c)), tf_ctl_na_rbf(__fmul_rn(r, s)));
+        }
+        row[i] = __float2bfloat16_rn(v);
+    }
+}
+
+// Gate 1: nemo_audio's relu as fmaxf(x, 0), which turns a NaN into 0 where tf_na_relu keeps it.
+extern "C" __global__ void __launch_bounds__(256) tf_ctl_na_relu_fmax(__nv_bfloat16* x, int64_t count) {
+    const int64_t at = static_cast<int64_t>(blockIdx.x) * 256 + threadIdx.x;
+    if (at < count) x[at] = __float2bfloat16_rn(fmaxf(__bfloat162float(x[at]), 0.0f));
+}