drama: the backbone chain launches with programmatic dependent launch by plotarmordev · Pull Request #927 · ashhart/TensorFold

plotarmordev · 2026-10-16T08:13:39Z

What this changes

The backbone's kernels use programmatic dependent launch and wait before accessing predecessor-produced data; the
one-row GEMV may prefetch immutable weights before its wait (part of #914; stacked on #926, its own commit is
the last one). PDL is enabled only for embedded backbone images that both carry the drama_pdl marker read by
Module.holds (#917); standalone kernel directories use plain launches.

This stack implements #914 and depends on #915 (allocation cap), #916 (direct-read fallback) and #917
(Module.holds), inherited through #920.

This inherits the stack's overlap with #844, including the recorded CUDA integration and CHANGELOG.md conflicts;
land #844 first, then rebase the stack.

Receipt

Environment, with the TensorFold commit, MLX or PyTorch version, machine and GPU, checkpoint and revision:
- Code. Head bef432bbb25e, on main c7137bb3. No MLX or PyTorch.
- Builds and host tests. An NVIDIA GB10 build host on Ubuntu 24.04.4 LTS aarch64 (Linux 6.17), with Zig
  0.17.0, nvcc 13.0.88 and CUDA_VISIBLE_DEVICES empty.
- GPU measurements. A GB10 (sm_121) host with CUDA 13 libraries (cuBLASLt 13.1.1). Checkpoint
  fishaudio/drama-3 at e0a9218a; reference sglang-omni ab095ef9.
- Which build was measured. These commits built on TensorFold 1e76a70e. Here they are rebased onto
  c7137bb3 with comments made single-line and files split. Every Drama fatbin of this head is byte-identical to
  the measured series' at the same commit. The host Zig is token-identical apart from the file split, compared
  with std.zig.Ast.
Exactness, from tools/bench_concurrent.py --alone --serial or token_sha against "draft": false: n/a for
tokens. Each speed stage was gated on all 46 normal replies of a 48-request corpus capture: byte-identical to the
build before any speed change. The two capped long-05 replies were identical to the
reviewed clipped capture, and bench/stream_check.py passed 9/9 (speed-stages.json). Fidelity against the reference is not qualified
(qualification.json); preserving native output does not establish reference equivalence.
Decode speed before and after, from tools/bench_openai.py: n/a, because it measures tokens. Warm real-time factor
(speed-stages.json):
- 0.678 → 0.655 for stage 05, this commit;
- the promoted final build: 0.636;
- on the final build against the reference, 0.634 against 0.939 (before-after-02). The reference/final
  comparison is a separate benchmark from the stage timings, with differing audio durations and cache/admission
  conditions whose timing effect was not isolated (qualification.json).
Prompt speed before and after: the cold first request's ttft_s from tools/bench_cache.py --require-cold on a long prompt, or "n/a" with the reason when the change cannot reach prompt processing: Not measured. Prompt kernels participate in the changed launch chain. Final-series
first audio is 0.141 s p50 (before-after-02), but that is not an isolated prompt-processing before/after
measurement.
Tests run, with pass, fail and skip counts:
- GPU audio gate on the stack's tip bef432bbb25e (one GB10, under the recipe's guard): all 46 normal-request
  audio bodies byte-identical to the promoted pre-upstream build, the long-05 clipped body identical in both
  passes, bench/stream_check.py 9/9, warm RTF 0.654 against the pre-speed 1.087, and all six tower one-row
  projections on drama_tower_gemv.
- CUDA_VISIBLE_DEVICES= zig build test-cuda-host --summary all from an empty local cache: 29/29 steps, 498 passed, 0 failed, 0 skipped.
- zig build native -Dnvcc=/usr/local/cuda/bin/nvcc --summary all: 47/47 steps.
- zig build fatbins -Dnvcc=/usr/local/cuda/bin/nvcc: every Drama fatbin is byte-identical to the measured
  series' at this commit.
- python3 tools/zig/lean_check.py: 527 problems at every commit, main's 527 (lean_new.py: 0 new).
- zig fmt --check: clean at every commit.
- zig build test-golden --summary all: 4/4 steps.
- zig build STEP -Dnvcc=/usr/local/cuda/bin/nvcc --summary all for each drama-* tool step of this head:
  19/19 steps build.
Not run, and why:
- tf-cuda-test drama-pdl checks the PDL decision for the embedded images and a kernel directory. It needs a
  GPU and was not run for this head.
- A GPU run on this exact head: the GPU audio gate ran once, on the stack's tip bef432bbb25e; this head is
  host-tested only.
- Metal: the family is CUDA-only.
- tools/bench_openai.py, bench_concurrent.py and bench_cache.py: they drive text tokens, so speech used the
  recipe's harness.

Checklist

Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off. (Drama 3
generates audio frames in its own host loop, as in drama: Drama 3 (fishaudio/drama-3) text-to-speech on CUDA #920. Text families are unchanged.)
Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (No
drafts or resume. The measured speed stages preserved the 46 normal corpus audio bodies; concurrent-equals-solo
is not measured.)
No precision traded for speed. If the bits change, the description says which sums change and why. (Launch order and dependencies only; the measured speed stages preserved the 46 normal corpus audio bodies.)
Prompt processing is no slower than the last release. (Not measured. Text prompt paths are not touched.)
The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
New tests fail before the change, pass after it, and skip cleanly without their dependency. (tf-cuda-test drama-pdl is a GPU test and was not run here.)
Comments and docstrings are one line. No measurements or history in the source.
No personal data, machine names, internal hosts or local paths. No AI attribution lines.

setAllocationLimit bounds the bytes of device buffers plus pinned host buffers. An allocation reserves its bytes before the driver call and a failed call refunds them, so concurrent allocations cannot overshoot the cap and a refused one leaves usage and the measured peak unchanged. Host counts now take the same lock as device counts. An engine that admits a model against a memory budget can hold its own allocations to it.

The tower turns a backbone row into a frame's ten codebook values: four shared and nine residual layers whose cache keeps every preceding text and audio row, a semantic head and the reference sampler (temperature, top-p over the whole vocabulary, then top-k). Its projections run on cuBLASLt; drama_tower.cu holds the mixer, norms, RoPE, attention and sampling kernels.

The codec decodes ten codebooks a frame to mono 44.1 kHz PCM, 2048 samples a frame, as the reference's FP16 vocoder does: RVQ, the windowed post-module transformer, two upsample stages and the DAC decoder. codec_pth.zig reads torch.save's zip archive and runs its state-dict pickle on a VM that knows only the opcodes and globals a tensor state dict uses, so nothing in the file runs. Weight norm is folded in FP32 and tensors are cast to FP16 with round-to-nearest-even, as the reference does. codec.zig holds the decode; codec_load.zig holds the codec's geometry, the weight-norm folding, the FP16 casts and the one device weight image.

An engine may now synthesize speech: SpeechRequest carries a prompt, sampling settings and caller-owned PCM, Info.speech describes the audio it produces, and Info.text_generation lets a speech-only engine refuse the text routes. Text submit gains error.Unsupported, which the chat route reports as a request error.

The native CUDA backend opens fish_omni checkpoints with the Drama 3 host: one synthesis at a time on a worker thread that owns the context, stream, backbone, tower and codec, admitted against the same memory pool the text families use. The pool reading moves to cuda_memory.readPool so both paths share it, and the host caps its own allocations at what it was admitted for. Options the host cannot honor are refused at open.

OpenAI's speech request, built into Drama 3's default-voice prompt as the reference's prompt builder builds it and synthesized whole, then sent as WAV or raw PCM. Fields the route would silently drop are refused, and a speech-only engine starts the server without the text-only services.

Drivers that check the native backbone and tower against captured reference tensors, replay single operators against them, and trace each stage for attribution, plus tf-cuda-test drama-conv and drama-qk on synthetic rows. None of them runs as part of the host tests.

The codec can decode frames from a given frame onward at their absolute rows, reading the rows before them that its causal convolutions and attention windows reach, and writing nothing before them. Every chunk boundary gives the same samples as one decode of all frames, so audio can be decoded while it is generated.

Drivers that decode captured codes with the native codec for comparison with the reference decoder, replay its RVQ and weight preparation in isolation, and trace each decoder stage. None of them runs as part of the host tests.

… generates A streamed speech request now receives audio events: each one hands the caller a final, rising prefix of the PCM before the finished event. The Drama host decodes the first frame alone, for the earliest audio, then frames in groups, and publishes each decoded prefix.

…ve host and their tool steps Embeds the backbone and tower fatbins, adds the drama module to the native server and its host tests, and a build step per backbone and tower oracle, replay and trace driver. The codec's kernels and tools follow on their own.

… are decoded With stream: true and response_format pcm, the route sends HTTP/1.1 chunks as the engine publishes audio. A stream that fails after its first chunk ends without a terminator, so a client cannot mistake it for complete audio.

…memory A single-row projection copies each warp's weight chunks into a shared ring with cp.async ahead of use and reads them back with ldmatrix, issuing the same mma sequence on the same fragments in the same K order as the row tiling it replaces, so the results are the same bits.

Embeds the codec fatbin and adds a build step per codec oracle, replay and trace driver. With this the series builds.

Decodes the same frame codes whole and in the host's chunk policy, one frame at a time and in irregular chunks, each over NaN-poisoned scratch, and requires every chunked decode to equal the whole decode.

Some file systems accept an O_DIRECT open and then fail the aligned read with EINVAL. The reader now clears O_DIRECT on that descriptor and reads again through the page cache instead of failing the load; a test injects the refusal.

A marker test: an image can carry a device global that promises something about its kernels, and the host can check for it before relying on the promise. Any driver failure, a missing name included, answers false rather than an error. A test drives it through a fake driver table.

Drama 3's backbone is a Qwen3.5-35B-A3B-style hybrid: 30 Gated DeltaNet and 10 gated full-attention layers, each with a 256-expert top-8 block-FP8 MoE and a gated shared expert. Model.run appends rows and returns every row's final-norm hidden state and the last row's semantic logits. The kernels follow the arithmetic of the sglang-omni reference: the canonical 64-token chunked GDN prefill (drama_gdn_chunk.cuh, after flash-linear-attention's kernels as SGLang carries them), the packed recurrent decode, gated attention with Q/K norm and partial NeoX RoPE, dynamic 1 x 128 FP8 activation quantization, and the BF16 tensor-core GEMMs in drama_backbone_gemm.cu. Weights load as stored; nothing is requantized. THIRD_PARTY_NOTICES.md credits flash-linear-attention (MIT), whose chunked kernels the GDN prefill is adapted from and whose packed decode the recurrence follows.

…ync ring The one-row path keeps the same mma sequence and fragments as the tilings it replaces, now for the semantic head's grouping too, while its weights stream through an unpadded shared ring whose 16-byte pieces are swizzled so each read still hits distinct banks.

…ation The residual add and RMS norm before the MoE also write the normed rows' FP8 groups, in the same order and rounding as the separate quantization launch they replace.

The routed rows are stored as their FP8 quantization and the shared expert's rows as BF16, in one launch instead of two, with the same arithmetic per row.

Each thread of drama_moe_combine now writes 8 contiguous outputs from 16-byte loads; every sum keeps its order.

A warp loads and scores its keys in batches with the next batch's loads issued early; the running softmax update stays in key order, so the result is unchanged.

A cluster of CTAs shares a head: each scores its share of the key tiles, all gather the scores over distributed shared memory and replay the tile updates in order, then each sums its own slice of the head. Every value is computed as the one-CTA kernel computes it; prompt rows keep that kernel.

… order stager.zig collects each tensor's copy as the loaders walk their names, sorts them by file offset and streams each file through two pinned slots, reading one with O_DIRECT while copying from the other, in pieces that never span two device allocations.

…udio A synthesis that reaches max_frames without an end of speech keeps the audio of those frames, as the reference returns its capped replies, instead of failing without audio. A whole reply carries X-Finish-Reason stop or length; a stream declares the trailer and ends with X-Finish-Reason: length only when it was clipped.

…ame launch For a single row, the attention launch also does drama_attn_prep's norm, RoPE and cache append, taking the new key and value from registers, with the same roundings as the separate launch.

The sampler's ranking comes from a stable LSD radix sort under the same strict total order the bitonic network used, so the permutation and every sum over it are unchanged.

…one loads The codec's host-side weight image is CPU work that the backbone and tower reads leave room for; it now runs on its own thread with the context current there, and the host joins it before it finishes opening.

…Lt's summation order drama_tower_gemv sums in the order of the gemv kernel cuBLASLt picks for the tower's one-row plans. At init each projection is checked against its cuBLASLt plan bit for bit and uses the kernel only if they agree. The qualification and its probe weights live in tower_gemv.zig.

…uster guard

…m_90 cluster guard zig/src/families/drama/tower.zig: the guard beside 926's pub projections and loadModule.

…he one-CTA kernel below it

…n types in their own files, Drama's refusals in cuda_refusal Upstream main's engine_api.zig is at 600 lines and native/cuda.zig at 583, so the merge took them to 632 and 602. The speech request, result, offer and submit error move to core/speech_api.zig, and the image and video input types to core/vision_api.zig, re-exported under their names as rows_ask.zig's are. Drama 3's flag refusals join LFM2's in cuda_refusal.zig, with a test. No declaration, behavior or test is removed.

…ion types in their own files zig/src/core/engine_api.zig: the speech types leave as on 920; the streamed request's comment and stream field move with them into core/speech_api.zig.

…d vision types in their own files

…and vision types in their own files

…vision types in their own files

…ypes in their own files

…sion types in their own files

…pes in their own files

ashhart#921 adds Event.audio; main's newer lane host tests switch over every event, so zig build test stopped compiling them.

plotarmordev · 2026-10-17T06:07:58Z

Thank you, @ashhart. The combined family plan works for us.

sm_89 cluster path. drama: faster tower decode attention and sampling, same bits #924 now carries the guard: images built below sm_90 leave drama_attention_decode out, and the tower runs drama_attention for the decode row instead.
Controls stay visible. We keep the component-fidelity failures, the long-text cases and the whole/chunk controls in the series as they are. Each optimization will be re-proved against the merged baseline once it exists, rather than argued from earlier native audio hashes.
CI. zig build test fails on speech: streamed Drama 3 audio on /v1/audio/speech #921-drama: the backbone chain launches with programmatic dependent launch #927 because three lane tests from main's newer commits don't handle the new Event.audio. A fix is coming on speech: streamed Drama 3 audio on /v1/audio/speech #921, merged forward through the stack.
A fidelity fix to offer. Full attention now runs the reference's 64-key softmax tiles, the same tiling as the reference. On our GB10, against sglang-omni ab095ef9 on a 182-token text fixture:
- worst-row cosine 0.8006 → 0.9169 (bar 0.8713)
- semantic KL 0.0243 → 0.0062 (bar 0.01)
- top-30 overlap 0.867 → 0.933 (bar 0.9)
- silent stalls 2 → 0 on our 40 test replies, but only with our own test voice and long-form settings. On upstream's default voice with no flags we see no improvement, so we don't claim one.
- about 5% slower
All-row cosine, last-row cosine and relative L2 still miss their bars narrowly, so this is testing, not qualified. It will come as a separate PR stacked on this series, with its component cells and corpus run on our hardware, for you to take into the combined family or leave.

…s ignore the speech audio event

…sts ignore the speech audio event

…ignore the speech audio event

…he speech audio event

…nore the speech audio event

…e speech audio event

ashhart · 2026-10-17T22:12:35Z

Thanks @plotarmordev for the sm_89 guard and for keeping the fidelity failures visible. The combined family remains the plan, with both contributor chains credited and each speed increment re-proved against that baseline.

#985 is useful testing work, but the improved 0.92 worst row and 0.006 KL do not close the remaining all-row, last-row and relative-L2 bars. The default-voice result also stays explicit. We will not describe it as qualified or lower those bars to pass it. Keep the event-test repair, chunk controls and complete matched audio checks with the final integrated family.

Posted by Ash's AI Assistant Synthena

…tal, with #915-#926 and #939 beneath them #927's stack brings the process-wide cap on device and pinned host bytes, direct reads a file system refuses going through the page cache, Module.holds, the Drama 3 backbone, tower and speech codec on CUDA (codec.pth read without Python), speech requests and audio events in engine_api, POST /v1/audio/speech with streamed PCM, and its decode kernels; #956's (with #939) the family's shared parts, its Metal backend and drama3-generate on macOS. Laid as the maintained series: the bridges between the two, qualification, and owners that keep buffers and modules until their completion is proved. On this line the cap's reservations join the pinned peaks' accounting, the learner's API moves to core/learning_api.zig with engine_api.zig re-exporting it (Example keeps its context row), the CUDA family list holds LFM2 and fish_omni, and both build files keep this line's layout with Drama's steps added. The tests are named for what they check.

tools/drama3/oracle_prompt.py is a copy of SGLang-Omni's prompt.py, and the oracle, its model half and the codec converter import it; code copied from another project does not land, so the four scripts stay on #927's and #956's branches with the 25 archived #927 diagnostic programs (oracles_927/*.zig.txt, whose originals are in #927's merged history) and the series' provenance note. Serving reads the checkpoint and codec.pth in Zig, so nothing a user runs changes; the notices and the codec's comments no longer name the scripts.

… under Unreleased Bullets for the Qwen 35B and its first-token path, Bolt (with its decode as measured beside #897's own build), Omni video, Gemma 4, Step-5, H3 on Metal and on CUDA, Flash Next's shared sessions on Metal, ModelOpt below sm_90, the digest cache, Drama 3, Qwen-Image-2.1 and the two benchmark fixes. #927's own Unreleased block had landed between 1.0.12 and 1.0.11; its four lines are one bullet in the right section now, its timings given as its authors' own.

ashhart · 2026-10-18T10:55:37Z

Thank you, @plotarmordev, for the speech protocol and infrastructure, and @drowzeys for the family, backends and codec. 1.1.0 consolidates #915, #917, #920, #921, #922, #923, #924, #925, #926, #939 and #956 through #927 (5dcf77d7c). One CUDA/Metal Drama 3 family uses bounded shared frame rounds, request-owned state, native checkpoint/codec loading and streamed PCM speech. Tensor/frame bounds, staging and completion ownership are checked. HTTP serves the checked default-voice subset; experimental attention and codec variants remain separate. Audio fidelity, request interleaving and completion results: Apple M5 Ultra norm checks passed 1/1 and CUDA resources passed; checkpoint audio fidelity, seeded concurrency and cancellation remain unrun.

Posted by Ash's AI Assistant Synthena

plotarmordev added 30 commits October 16, 2026 12:18

tests: Drama 3 codec oracle, replays and stage trace

93fcaab

Drivers that decode captured codes with the native codec for comparison with the reference decoder, replay its RVQ and weight preparation in isolation, and trace each decoder stage. None of them runs as part of the host tests.

build: the Drama 3 codec kernels and the codec tool steps

10796d4

Embeds the codec fatbin and adds a build step per codec oracle, replay and trace driver. With this the series builds.

tests: tf-drama-codec-stream-oracle, whole against chunked codec decodes

8bc6814

Decodes the same frame codes whole and in the host's chunk policy, one frame at a time and in irregular chunks, each over NaN-poisoned scratch, and requires every chunked decode to equal the whole decode.

drama: add-norm fused with the routed experts' FP8 activation quantiz…

9dc8202

…ation The residual add and RMS norm before the MoE also write the normed rows' FP8 groups, in the same order and rounding as the separate quantization launch they replace.

drama: one SwiGLU launch for the routed and shared experts

705541c

The routed rows are stored as their FP8 quantization and the shared expert's rows as BF16, in one launch instead of two, with the same arithmetic per row.

drama: MoE combine on 8-wide vectors

08291b5

Each thread of drama_moe_combine now writes 8 contiguous outputs from 16-byte loads; every sum keeps its order.

drama: decode attention keeps the next keys' loads in flight

641b18b

A warp loads and scores its keys in batches with the next batch's loads issued early; the running softmax update stays in key order, so the result is unchanged.

drama: one-row attention runs its Q/K norm and RoPE prologue in the s…

9b3934c

…ame launch For a single row, the attention launch also does drama_attn_prep's norm, RoPE and cache append, taking the new key and value from registers, with the same roundings as the separate launch.

drama: the tower samples with a stable radix sort

ea06b8b

The sampler's ranking comes from a stable LSD radix sort under the same strict total order the bitonic network used, so the permutation and every sum over it are unchanged.

plotarmordev added 12 commits October 17, 2026 10:33

Merge align/926-merge-main into perf/drama3-pdl: the tower's sm_90 cl…

bc31ec2

…uster guard

Merge align/925-merge-main into perf/drama3-tower-gemv: the tower's s…

4d3c4f5

…m_90 cluster guard zig/src/families/drama/tower.zig: the guard beside 926's pub projections and loadModule.

CHANGELOG: the tower's cluster decode attention runs from sm_90 on, t…

2d62476

…he one-CTA kernel below it

Merge align/920-merge-main into feat/drama3-streaming: speech and vis…

f5aa7a5

…ion types in their own files zig/src/core/engine_api.zig: the speech types leave as on 920; the streamed request's comment and stream field move with them into core/speech_api.zig.

Merge align/921-merge-main into fix/drama3-clipped-replies: speech an…

fd87532

…d vision types in their own files

Merge align/922-merge-main into perf/drama3-backbone-one-row: speech …

572afe3

…and vision types in their own files

Merge align/923-merge-main into perf/drama3-tower-decode: speech and …

1ecf53b

…vision types in their own files

Merge align/924-merge-main into perf/drama3-load: speech and vision t…

bd11ba4

…ypes in their own files

Merge align/925-merge-main into perf/drama3-tower-gemv: speech and vi…

f495030

…sion types in their own files

Merge align/926-merge-main into perf/drama3-pdl: speech and vision ty…

f64f02d

…pes in their own files

core: lane tests ignore the speech audio event

118086a

ashhart#921 adds Event.audio; main's newer lane host tests switch over every event, so zig build test stopped compiling them.

plotarmordev added 6 commits October 17, 2026 14:10

Merge align/921-merge-main into fix/drama3-clipped-replies: lane test…

4d41f5a

…s ignore the speech audio event

Merge align/922-merge-main into perf/drama3-backbone-one-row: lane te…

2b27b19

…sts ignore the speech audio event

Merge align/923-merge-main into perf/drama3-tower-decode: lane tests …

e45d71b

…ignore the speech audio event

Merge align/924-merge-main into perf/drama3-load: lane tests ignore t…

809e103

…he speech audio event

Merge align/925-merge-main into perf/drama3-tower-gemv: lane tests ig…

0ed8946

…nore the speech audio event

Merge align/926-merge-main into perf/drama3-pdl: lane tests ignore th…

d986914

…e speech audio event

plotarmordev mentioned this pull request Oct 17, 2026

drama: backbone attention runs the reference's 64-key softmax tiles #985

Open

4 of 9 tasks

cwschroeder mentioned this pull request Oct 17, 2026

studio: Wasserman Slate AI Studio, an AI film studio on TensorFold (Mac + DGX Spark) #941

Open

This was referenced Oct 18, 2026

cuda: a process-wide cap on DeviceBuffer and pinned HostBuffer bytes #915

Open

cuda: Module.holds, whether a loaded image defines a device global #917

Open

This was referenced Oct 18, 2026

drama3: Fish Audio Drama 3 text to speech on one CUDA GPU (new family, plain decoding; proposal) #939

Open

drama3: Fish Audio Drama 3 text to speech on Metal, from Zig (stacked on #939) #956

Open

Navigation Menu

Repository navigation

drama: the backbone chain launches with programmatic dependent launch - #927

drama: the backbone chain launches with programmatic dependent launch#927
plotarmordev wants to merge 59 commits into
ashhart:mainfrom
plotarmordev:perf/drama3-pdl

plotarmordev commented Oct 16, 2026

Uh oh!

plotarmordev commented Oct 17, 2026 •

edited

Loading

Uh oh!

ashhart commented Oct 17, 2026

Uh oh!

ashhart commented Oct 18, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

plotarmordev commented Oct 16, 2026

What this changes

Receipt

Checklist

Uh oh!

plotarmordev commented Oct 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ashhart commented Oct 17, 2026

Uh oh!

ashhart commented Oct 18, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

plotarmordev commented Oct 17, 2026 •

edited

Loading