music3: MiniMax Music 3 lyrics to song on one CUDA GPU, from Zig (new family; proposal)

[drowzeys]

Thank you, Ash, for the invitation in #941, and for the CUDA runtime and kernel contract this builds on.

This is the first music model on TensorFold's Zig runtime. MiniMax Music 3 turns a caption and lyrics into a full song with vocals. It now runs end to end on TensorFold's CUDA runtime, with no Python and no PyTorch at run time. It reads Comfy-Org's single-file checkpoints as they are and has its own kernels.

The model has three parts:

• an 8B autoregressive LM with a depth decoder, which writes 8 audio codes a frame at 25 frames a second;
• a 2.4B flow-matching DiT;
• the DAV waveform decoder.

On one DGX Spark (GB10), for a 60 s song:

• 123.9 s with int8 weights, against 289 s for ComfyUI's native nodes on the same GPU. That is 2.3x faster.
• 17.9 GB peak, against 29.3 GB.
• The AR runs at its measured read ceiling: 20.4 frames/s int8 and 10.4 bf16, against ceilings of 20.7 and 10.4.
• Parity with ComfyUI at every stage: code-0 logits cos 0.999977, DiT final latent 37.6 dB, DAV 71 dB (bf16).
• Kernel contract: 16 of 16 kernels pass gates 1–4 (the 4 new music3 ones among them), 5 controls red as planted, 0 limit problems.
• Lean check: 0 new findings.

I did not open an issue first, since this came by invitation in #941. Happy to move the gate discussion into one if you prefer. It stands alone on main: it does not depend on the Drama 3 pull requests.

Receipts

All on one GB10 (sm_121, 128 GB), idle, one model process, CUDA 13. The checkpoint is Comfy-Org's MiniMax-Music3 repack (minimax_music3_text_encoder_pruned_bf16, minimax_music3_dit_fp16, minimax_music3_dav). The baseline and reference is ComfyUI 5a665890's native nodes on the same machine. The full tables and their JSON are in docs/recipes/minimax-music3-cuda.md and docs/recipes/minimax-music3-cuda/.

Measured ceilings before building

Measured on an idle GB10, with weights rotated through a 2 GB pool so L2 never holds them:

Measurement	Result
Plain read	230 GB/s
Warp-a-row GEMV, two activation rows, bf16 and int8	220–240 GB/s
cuBLAS fp16 GEMM, fp32 accumulation, at the DiT's shapes	96–104 TFLOP/s
AR ceiling from those reads	10.4 frames/s bf16, 20.7 frames/s int8
AR reached	10.4 frames/s bf16, 20.4 frames/s int8

What got there:

• CFG rows share weight reads: the conditional and unconditional rows go through each GEMV together.
• The whole frame step is one CUDA graph: the lm_head, the CFG top-k draw on the device, the depth decoder's 7 passes, and the frame fed back through 36 layers. The host reads the drawn codes every 16 frames for the stop code.
• PDL: each GEMV's first weight loads are in flight before griddepcontrol.wait.
• Fused SwiGLU GEMV: gate and up write the SwiGLU directly.
• Our own tensor-core flash attention for the DiT, with rope applied as q and k load (mma.sync m16n8k16). There is no flex_attention, which gives wrong results on sm_121.
• Timed cuBLASLt algorithms: each DiT GEMM shape's algorithm is timed once at load.
• Bias moved out of the GEMM epilogue, into the GLU and the next LayerNorm. The bias + accumulate epilogue picked 14 TFLOP/s kernels at some shapes.

Parity against ComfyUI

Same 30 s input (hiphop bench prompt, seed 7), teacher-forced so the rows are comparable:

Check	bf16	int8 (default)
Prompt ids	equal (and 10 of 10 id cases)	equal
Frame states cos, mean / min	0.999956 / 0.999556	0.999827 / 0.998345
Code-0 logits cos	0.999977	0.999897
Code-0 greedy agreement	97.5%	96.7%
ComfyUI's drawn code, median rank in ours	1	1
Depth logits cos, codebooks 1–7	0.999957–0.999966	0.99978–0.99984
DiT probes cos, sigma 1.0 / 0.5 / 0.1	≥ 0.999958	(DiT is fp16 in both)
DiT final latent (ComfyUI's noise and states)	cos 0.999914, 37.6 dB
DAV of ComfyUI's latent, chunked as served	cos 1.0000000, 71.0 dB

The draws themselves cannot match (torch's multinomial against our Philox inverse-CDF draw), so a song is not ComfyUI's song for the same seed, only the same model.

Speed and memory

Hiphop bench prompt, seed 7, and the long-pop prompt, seed 5, for the 120 s row. Load (37 s from a cold page cache) is outside "total". Peak machine is MemTotal − MemAvailable sampled twice a second, minus the idle machine.

Song	Weights	AR s (frames/s)	DiT s	DAV s	Total s	RTF	Peak machine GB
30 s	int8	36.4 (20.9)	22.0	1.9	60.3	2.01	16.6
60 s	int8	73.9 (20.4)	46.4	3.6	123.9	2.07	17.9
120 s	int8	151.7 (19.9)	93.0	7.3	252.0	2.10	19.3
30 s	bf16	72.1 (10.5)	21.9	1.9	96.0	3.20	24.0
60 s	bf16	144.4 (10.4)	46.3	3.6	194.3	3.24	25.5
120 s	bf16	292.7 (10.3)	93.1	7.3	393.1	3.28	26.6
60 s	ComfyUI baseline	176 (~9)	96	4	289	4.8	29.3

int8 is the default:

• It halves the bytes each frame reads.
• Its parity is within a few 1e-4 of bf16, and ComfyUI's drawn code is rank 1 in both.
• Its lyric accuracy is no lower than bf16's on three songs.

--weights bf16 reads the checkpoint's own numbers.

Lyric accuracy

Three 45 s songs, each in bf16 and int8, plus ComfyUI's own 60 s song. Transcribed by faster-whisper large-v3-turbo on the CPU and scored against the lyrics by tools/music3/lyrics_check.py. The score counts the words heard up to the last lyric word heard. The span counts only matched runs of 3 or more words.

Song	bf16 score / span	int8 score / span	Words heard, bf16 / int8
hiphop, seed 7	96 / 96	98 / 98	130 of 135 / 127 of 130
pop, seed 11	100 / 100	100 / 100	78 of 78 / 59 of 59
long-pop, seed 5	90 / 89	31 / 100	50 of 56 / 52 of 52
ComfyUI hiphop 60 s (reference)	96 / 96		159 of 166

The int8 long-pop score of 31 is a scoring artefact: a stray "down" matched the bridge's "go down". Its span is 100.

Checks on this branch

Run on a GB10 at this head, against main 6ee828c:

• zig build test: 86/86 steps; 450/456 tests passed, 6 skipped (libjpeg-turbo absent: JPEG cases).

• zig build test-cuda-host: 41/41 steps; 856/856 tests passed (main: 852/852).

• Plain zig build: succeeded.

• zig fmt --check zig build.zig: clean.

• python3 tools/zig/lean_check.py: 517 problems, the same 517 as main. 0 new findings; the inherited ones are unchanged.

• Kernel contract (zig build kernel-contract -Dnvcc=/usr/local/cuda/bin/nvcc) on sm_121:

  SUMMARY 16 kernels: 16 pass gates 1-4, 0 fail, 0 known findings; 5 controls: 5 red as planted, 0 missed; 0 limit problems
  

• Launch limits (zig build fatbins kernel-contract-fatbins -Dsm=80,86,89,120,121, then tf-cuda-test kernel-contract-limits zig-out/fatbin):

  PASS kernel-contract-limits: every registered launch fits sm_80, sm_86, sm_89, sm_120 and sm_121, every control red
  

  All 55 music3 launches are registered in registry.zig with records for sm_80, sm_86, sm_89, sm_120 and sm_121. m3_rmsnorm, m3_gemv_bf16_2, m3_gemv_i8_2 and m3_gemv_g8_2 go through gates 1–4 (zig/tests/cuda/contract_music3.zig).

What it adds

These are new files, plus the minimal wiring: 11 lines in zig/build/cuda.zig, 1 in zig/src/cuda/kernels.zig, 12 in the contract runner, 55 records in registry.zig, and 1 row in the recipe index.

• zig/src/families/music3/ (Zig over the CUDA runtime):
  • config.zig holds the shapes, request settings and the reference's schedules.
  • prompt.zig builds the caption and lyrics prompt ids through the core tokenizer and the checkpoint's own tokenizer_json.
  • weights.zig reads the three files and quantizes int8 on the GPU at load: body and lm_head one scale a row, the depth decoder one scale a group of 128.
  • ar.zig holds the AR prompt pass (cuBLAS bf16 GEMMs) and the graphed frame step.
  • dit.zig holds the windowed DiT: 689 latent rows every 344, both CFG branches in one forward, on cuBLASLt fp16 GEMMs.
  • dav.zig holds the DAV: fp32, every conv as one SGEMM a tap, with no TF32.
  • engine.zig and generate.zig provide the request, the CLI and --serve.
  • parity.zig and npy.zig provide the parity dumps.
  • lt.zig and blas.zig open cuBLASLt and cuBLAS at run time.
  • wav.zig writes the WAV.
• zig/kernels/cuda/music3.cu and its headers (music3_ar.cuh, music3_dit.cuh, music3_dav.cuh, music3_sample.cuh, music3_common.cuh) build one fatbin.
• zig build music3-generate:
  • caption and lyrics to a WAV;
  • --ids for prompt ids with no GPU;
  • --parity-ar / --parity-probe / --parity-dit / --parity-dav REF against a reference dump;
  • --serve 1.
• tools/music3/ (development tools, Python):
  • parity.py builds the parity table from REF and MINE, with numpy only.
  • lyrics_check.py measures lyric accuracy with Whisper.
  • served_test.py drives --serve through the Studio's client.
• docs/recipes/minimax-music3-cuda.md is the recipe, with the plan and ceilings written before building, the receipts, and their JSON.

How it fits the engine, and how it does not

• Serving. --serve 1 reads one JSON request a line on stdin and answers on stderr, with optional progress lines. It is the same protocol as our Mac MLX engine for this model, so one client drives either. It answers one request at a time.
• Lanes. As with Drama 3 (drama3: Fish Audio Drama 3 text to speech on one CUDA GPU (new family, plain decoding; proposal) #939), a step here is a frame of 8 codes, not one u32 token, and the song is finished by a DiT and a waveform decoder. So this runs as its own generator beside the lane engine rather than on the lane core. Moving it there needs the same core changes drama3: Fish Audio Drama 3 text to speech on one CUDA GPU (new family, plain decoding; proposal) #939 lists: a frame position type or a family hook, per-codebook sampling, and an audio output. Plain decoding only; there is no drafter.
• Native, but not yet one engine request. Nothing at run time needs Python, PyTorch or ComfyUI. ComfyUI is only the reference for the parity receipts.

Licences

• The code is written for TensorFold and is Apache-2.0. No ComfyUI source is copied in, and nothing here imports ComfyUI. ComfyUI's implementation (GPL-3) is the reference it was checked against. Some parts were written by hand to behave as ComfyUI does, so the outputs line up, and I am naming them so you can decide:

  • the prompt formatting rules;
  • the "simple" sigma schedule;
  • the fp16 rounding of the DiT's rope angles;
  • the top-k tie rule;
  • the DAV level rule.

• The reference harness is not included. It is a ComfyUI custom node that records tensors from ComfyUI's own loaders and sampler. The recipe describes what it records and the file layout parity.py reads.

• The model: MiniMax-Music3 is under the MiniMax-Music3 Community License. No weights are included.

  • A commercial product must display "MiniMax-Music3".
  • Use by a business above $20M a year in revenue needs MiniMax's written authorization.
  • A hosted service must keep the licence's safeguards.

  music3-generate prints the display notice, and --serve reports "license_display":"MiniMax-Music3".

Not tested

• GPUs other than GB10. The kernels build and pass gate 5 for sm_80/86/89/120, but nothing ran there. PDL is used only on sm_90 and up.
• fp8, int4 or NVFP4 weights, and fp8 DiT GEMMs.
• Songs over 120 s. The cap is 360 s and the caches are sized for it, but that length was not run.
• Prompts near 5,000 tokens, an instrumental-only lyric check, and non-ASCII section tags.
• Listening by ear beyond spot checks. Lyric accuracy is measured by Whisper.
• Two requests at once: the server answers one at a time.
• The kernel contract covers 4 of the 55 launches through gates 1–4. The rest have gate 5 records only.

Next steps, if you want them

• The other launches through gates 1–4, and the lane-core questions above, for the gate we agree on.
• A native Metal Music 3 family on the Zig runtime. Our MLX engine for this model (branch music3-mlx on our fork) is a working reference and baseline for it. On an M5 Ultra it does a 60 s song at about real time in int8, with the same parity tables against ComfyUI on MPS. It is Python/MLX, so I have not sent it here: CONTRIBUTING keeps the Python engine frozen at 0.6.6, and the final native request should not need Python.

[ashhart]

Thank you @drowzeys! Music 3 on CUDA and Metal is a big piece of work, and your parity receipts against the reference made it much easier to review. We've folded #982 and #987 into one family on the shared core: frames go through the core's frame queue, each lane keeps its own AR state, prompts resume from the core prompt store with both CFG histories, and every load step unwinds only after its GPU work completes.

One change: prompt.zig said it follows ComfyUI's prompt.py by hand. ComfyUI is another project's GPL-3.0 code, so we rewrote the prompt builder from its behaviour. It matches yours byte for byte on 2,066 cases.

Could you attach the ten prompt cases from your receipts, with ComfyUI's ids, as JSON? They're the last check before the GPU gates.

Posted by Ash's AI Assistant Synthena