retroSoC/scripts/generate_apu_p5_microcode.py

#!/usr/bin/env python3
"""Generate the reviewable APU-P5 WAV/FLAC release microassembly."""

from __future__ import annotations

import argparse
from pathlib import Path


class Program:
    """Small deterministic emitter for the frozen text microassembly syntax."""

    def __init__(self) -> None:
        self.lines: list[str] = []

    def line(self, value: str = "") -> None:
        self.lines.append(value)

    def label(self, name: str) -> None:
        self.lines.append(f"{name}:")

    def op(self, operation: str) -> None:
        self.lines.append(f"  {operation}")

    def text(self) -> str:
        return "\n".join(self.lines) + "\n"


CTX_CONTROL = 0x47C0
CTX_INPUT_CONFIG = 0x47C4
CTX_OUTPUT_CONFIG = 0x47C8
CTX_FLAGS = 0x47CC
CTX_PARSED = 0x47D0
CTX_DIAGNOSTIC = 0x47D4
CTX_SOURCE_INFO = 0x47D8
CTX_FINAL = 0x47DC
CTX_INPUT_LENGTH = 0x47E0
CTX_OUTPUT_CAPACITY = 0x47E4
CTX_OUTPUT_CURSOR = 0x47EC
READER_STAGING = 0x4380

WAV_STATE = {
    "warnings": 0x4400,
    "strict": 0x4404,
    "input_length": 0x4408,
    "riff_end": 0x440C,
    "fmt_seen": 0x4410,
    "data_seen": 0x4414,
    "channels": 0x4418,
    "rate": 0x441C,
    "bits": 0x4420,
    "align": 0x4424,
    "output_rate": 0x4428,
    "output_channels": 0x442C,
    "output_bits": 0x4430,
    "output_mode": 0x4434,
    "data_remaining": 0x4438,
    "chunk_count": 0x443C,
    "skipped": 0x4440,
    "pending": 0x4444,
    "pending_bytes": 0x4448,
    "pending_frames": 0x444C,
    "pending_input": 0x4450,
    "tile_frames": 0x4454,
    "downmix": 0x4458,
    "resample": 0x445C,
    "chunk_size": 0x4460,
    "chunk_pad": 0x4464,
    "profile": 0x4468,
    "queued_frames": 0x446C,
    "pack_input": 0x4470,
    "uncommitted_input": 0x4474,
    "source_tile_frames": 0x4478,
    "frame_pointer": 0x447C,
    "format": 0x4480,
    "pack_capacity": 0x4534,
    "average_rate": 0x4538,
    "pack_output": 0x4550,
    "pending_output": 0x4554,
    "normalize_dst": 0x4558,
    "resample_pending": 0x455C,
    "resample_final": 0x4560,
    "resample_queued": 0x4564,
    "resample_output": 0x4568,
    "capacity_stop": 0x456C,
    "error_code": 0x4700,
    "error_reason": 0x4704,
    "inflight": 0x4730,
    "inflight_frames": 0x4734,
    "inflight_input": 0x4738,
    "inflight_final": 0x473C,
}

FLAC_STATE = {
    "min_block": 0x4484,
    "max_block": 0x4488,
    "min_frame": 0x448C,
    "max_frame": 0x4490,
    "total_lo": 0x4494,
    "total_hi": 0x4498,
    "decoded_lo": 0x449C,
    "decoded_hi": 0x44A0,
    "metadata_done": 0x44A4,
    "metadata_count": 0x44A8,
    "metadata_bytes": 0x44AC,
    "block_size": 0x44B0,
    "assignment": 0x44B4,
    "strategy": 0x44B8,
    "strategy_seen": 0x44BC,
    "frame_number": 0x44C0,
    "frame_start": 0x44C4,
    "channel": 0x44C8,
    "output_base": 0x44CC,
    "kind": 0x44D0,
    "order": 0x44D4,
    "wasted0": 0x44D8,
    "wasted1": 0x44DC,
    "coded_bits": 0x44E0,
    "partition_order": 0x44E4,
    "partition_index": 0x44E8,
    "parameter_bits": 0x44EC,
    "rice_parameter": 0x44F0,
    "raw_width": 0x44F4,
    "residual_buffered": 0x44F8,
    "reconstruct_output": 0x44FC,
    "unary_done": 0x453C,
    "coefficient_pointer": 0x4548,
    "process_index": 0x454C,
    "frame_input": 0x4520,
    "frame_remaining": 0x4524,
    "parsing_stage": 0x4528,
    "expected_number": 0x452C,
    "parsed_number": 0x4530,
    "signed_width": 0x4708,
}


def _store(program: Program, register: int, offset: int) -> None:
    program.op(f"st32 r0, r{register}, 0x{offset:04x}")


def _load(program: Program, register: int, offset: int) -> None:
    program.op(f"ld32 r{register}, r0, 0x{offset:04x}")


def _fail(program: Program, code: int, reason: int, label: str = "wav_fail") -> None:
    program.op(f"movi r4, {code}")
    program.op(f"movi r5, 0x{reason:04x}")
    program.op(f"jump_fwd {label}")


def _compare_fail(
    program: Program,
    left: int,
    right: int,
    predicate: str,
    code: int,
    reason: int,
) -> None:
    ok = f"wav_check_ok_{len(program.lines)}"
    program.op(f"cmp r{left}, r{right}")
    inverse = {
        "eq": "ne",
        "ne": "eq",
        "ult": "uge",
        "uge": "ult",
        "slt": "sge",
        "sge": "slt",
    }[predicate]
    program.op(f"jump_fwd {ok} if={inverse}")
    _fail(program, code, reason)
    program.label(ok)


def _emit_wav(program: Program) -> None:
    # r9..r15 are the byte-reader state: bit byte/count, parsed count,
    # FIFO word/valid bytes, refill bytes, and return value.
    program.label("wav")
    program.op("movi r0, 0")
    program.op("movi r1, 0x4400")
    program.op("movi r2, 208")
    program.op("loop_setup 0, r2")
    program.label("wav_state_clear_loop")
    program.op("st32 r1, r0, 0")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.op("loop_back 0, wav_state_clear_loop")
    _load(program, 1, CTX_FLAGS)
    program.op("movi r2, 1")
    program.op("and r1, r1, r2")
    _store(program, 1, WAV_STATE["strict"])
    _load(program, 1, CTX_INPUT_LENGTH)
    _store(program, 1, WAV_STATE["input_length"])

    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("call_fwd wav_get_u32le")
    program.op("movi r1, 0x46464952")
    _compare_fail(program, 15, 1, "ne", 4, 0x0010)
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("call_fwd wav_get_u32le")
    program.op("movi r1, 8")
    program.op("add r2, r15, r1")
    program.op("cmp r2, r15")
    program.op("jump_fwd wav_bad_extent if=ult")
    program.op("movi r1, 12")
    program.op("cmp r2, r1")
    program.op("jump_fwd wav_bad_extent if=ult")
    _load(program, 1, WAV_STATE["input_length"])
    program.op("cmp r1, r2")
    program.op("jump_fwd wav_bad_extent if=ult")
    program.label("wav_riff_bound_ok")
    _store(program, 2, WAV_STATE["riff_end"])
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("call_fwd wav_get_u32le")
    program.op("movi r1, 0x45564157")
    _compare_fail(program, 15, 1, "ne", 4, 0x0010)

    # Each chunk consumes at least one byte; cap the frozen 1024-chunk bound
    # by the input length so short files do not spin through idle iterations.
    _load(program, 1, WAV_STATE["input_length"])
    program.op("movi r2, 1024")
    program.op("min r1, r1, r2")
    program.op("loop_setup 0, r1")
    program.label("wav_chunk_loop")
    _load(program, 1, WAV_STATE["riff_end"])
    program.op("cmp r11, r1")
    program.op("jump_fwd wav_chunk_idle if=eq")
    program.op("sub r2, r1, r11")
    program.op("movi r3, 8")
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_truncated if=ult")
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("call_fwd wav_get_u32le")
    program.op("mov r4, r15")
    program.op("call_fwd wav_get_u32le")
    program.op("mov r5, r15")
    _store(program, 5, WAV_STATE["chunk_size"])
    program.op("add r6, r11, r5")
    program.op("cmp r6, r11")
    program.op("jump_fwd wav_bad_chunk if=ult")
    program.op("movi r7, 1")
    program.op("and r7, r5, r7")
    _store(program, 7, WAV_STATE["chunk_pad"])
    program.op("add r8, r6, r7")
    program.op("cmp r8, r6")
    program.op("jump_fwd wav_bad_chunk if=ult")
    _load(program, 1, WAV_STATE["riff_end"])
    program.op("cmp r1, r8")
    program.op("jump_fwd wav_bad_chunk if=ult")
    program.label("wav_chunk_bound_ok")
    _load(program, 1, WAV_STATE["chunk_count"])
    program.op("movi r2, 1")
    program.op("add r1, r1, r2")
    _store(program, 1, WAV_STATE["chunk_count"])
    program.op("movi r1, 0x20746d66")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_fmt_chunk if=eq")
    program.op("movi r1, 0x61746164")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_data_chunk if=eq")
    _load(program, 1, WAV_STATE["skipped"])
    _load(program, 2, WAV_STATE["chunk_size"])
    _load(program, 3, WAV_STATE["chunk_pad"])
    program.op("add r2, r2, r3")
    program.op("add r1, r1, r2")
    program.op("cmp r1, r2")
    program.op("jump_fwd wav_bad_chunk if=ult")
    program.op("movi r3, 0x00100001")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_profile_limit if=uge")
    _store(program, 1, WAV_STATE["skipped"])
    _load(program, 1, WAV_STATE["chunk_size"])
    program.op("call_fwd wav_skip_bytes")
    program.op("jump_fwd wav_chunk_padding")

    program.label("wav_fmt_chunk")
    _load(program, 1, WAV_STATE["fmt_seen"])
    _load(program, 2, WAV_STATE["data_seen"])
    program.op("or r1, r1, r2")
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_required_chunk if=ne")
    program.op("movi r1, 16")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_fmt_size_ok if=eq")
    program.op("movi r1, 18")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_fmt_size_ok if=eq")
    program.op("movi r1, 40")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_fmt_size_ok if=eq")
    _fail(program, 4, 0x0014)
    program.label("wav_fmt_size_ok")
    program.op("call_fwd wav_get_u16le")
    program.op("mov r4, r15")
    program.op("call_fwd wav_get_u16le")
    program.op("mov r3, r15")
    _store(program, 3, WAV_STATE["channels"])
    program.op("call_fwd wav_get_u32le")
    _store(program, 15, WAV_STATE["rate"])
    program.op("call_fwd wav_get_u32le")
    program.op("mov r2, r15")
    _store(program, 2, WAV_STATE["average_rate"])
    program.op("call_fwd wav_get_u16le")
    program.op("mov r6, r15")
    _store(program, 6, WAV_STATE["align"])
    program.op("call_fwd wav_get_u16le")
    program.op("mov r7, r15")
    _store(program, 7, WAV_STATE["bits"])
    program.op("movi r1, 1")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_pcm_fmt if=eq")
    program.op("movi r1, 0xfffe")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_extensible_fmt if=eq")
    _fail(program, 3, 0x0003)

    program.label("wav_pcm_fmt")
    program.op("movi r1, 40")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_bad_format if=eq")
    program.op("movi r1, 18")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_pcm_cbsize if=eq")
    program.op("jump_fwd wav_validate_fmt")
    program.label("wav_pcm_cbsize")
    program.op("call_fwd wav_get_u16le")
    program.op("cmp r15, r0")
    program.op("jump_fwd wav_bad_format if=ne")
    program.op("jump_fwd wav_validate_fmt")

    program.label("wav_extensible_fmt")
    program.op("movi r1, 40")
    program.op("cmp r5, r1")
    program.op("jump_fwd wav_unsupported_encoding if=ne")
    program.op("call_fwd wav_get_u16le")
    program.op("movi r1, 22")
    program.op("cmp r15, r1")
    program.op("jump_fwd wav_unsupported_encoding if=ne")
    program.op("call_fwd wav_get_u16le")
    program.op("cmp r15, r7")
    program.op("jump_fwd wav_unsupported_encoding if=ne")
    program.op("call_fwd wav_get_u32le")
    program.op("mov r4, r15")
    program.op("movi r1, 1")
    program.op("cmp r3, r1")
    program.op("jump_fwd wav_extensible_mono if=eq")
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_extensible_guid if=eq")
    program.op("movi r1, 3")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_unsupported_encoding if=ne")
    program.op("jump_fwd wav_extensible_guid")
    program.label("wav_extensible_mono")
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_extensible_guid if=eq")
    program.op("movi r1, 4")
    program.op("cmp r4, r1")
    program.op("jump_fwd wav_unsupported_encoding if=ne")
    program.label("wav_extensible_guid")
    for expected in (0x00000001, 0x00100000, 0xAA000080, 0x719B3800):
        program.op("call_fwd wav_get_u32le")
        program.op(f"movi r1, 0x{expected:08x}")
        program.op("cmp r15, r1")
        program.op("jump_fwd wav_unsupported_encoding if=ne")

    program.label("wav_validate_fmt")
    program.op("movi r1, 1")
    program.op("cmp r3, r1")
    program.op("jump_fwd wav_channels_ok if=eq")
    program.op("movi r1, 2")
    program.op("cmp r3, r1")
    program.op("jump_fwd wav_bad_format if=ne")
    program.label("wav_channels_ok")
    _load(program, 1, WAV_STATE["rate"])
    program.op("movi r4, 8000")
    program.op("cmp r1, r4")
    program.op("jump_fwd wav_bad_format if=ult")
    program.op("movi r4, 96001")
    program.op("cmp r1, r4")
    program.op("jump_fwd wav_bad_format if=uge")
    program.op("movi r1, 8")
    program.op("cmp r7, r1")
    program.op("jump_fwd wav_bits_divide if=eq")
    program.op("movi r1, 16")
    program.op("cmp r7, r1")
    program.op("jump_fwd wav_bits_divide if=eq")
    program.op("movi r1, 24")
    program.op("cmp r7, r1")
    program.op("jump_fwd wav_bits_divide if=eq")
    program.op("movi r1, 32")
    program.op("cmp r7, r1")
    program.op("jump_fwd wav_bad_format if=ne")
    program.label("wav_bits_divide")
    program.op("movi r1, 3")
    program.op("shr r4, r7, r1")
    program.op("movi r1, 2")
    program.op("cmp r3, r1")
    program.op("movi r1, 1")
    program.op("shl r4, r4, r1 if=eq")
    _load(program, 6, WAV_STATE["align"])
    program.op("cmp r6, r4")
    program.op("jump_fwd wav_bad_format if=ne")
    _load(program, 1, WAV_STATE["rate"])
    program.op("movi r3, 1")
    program.op("cmp r4, r3")
    program.op("jump_fwd wav_average_ready if=eq")
    program.op("movi r3, 2")
    program.op("cmp r4, r3")
    program.op("movi r3, 1")
    program.op("shl r1, r1, r3 if=eq")
    program.op("jump_fwd wav_average_ready if=eq")
    program.op("movi r3, 3")
    program.op("cmp r4, r3")
    program.op("jump_fwd wav_average_three if=eq")
    program.op("movi r3, 4")
    program.op("cmp r4, r3")
    program.op("movi r3, 2")
    program.op("shl r1, r1, r3 if=eq")
    program.op("jump_fwd wav_average_ready if=eq")
    program.op("movi r3, 6")
    program.op("cmp r4, r3")
    program.op("jump_fwd wav_average_six if=eq")
    program.op("movi r3, 3")
    program.op("shl r1, r1, r3")
    program.op("jump_fwd wav_average_ready")
    program.label("wav_average_three")
    program.op("movi r3, 1")
    program.op("shl r2, r1, r3")
    program.op("add r1, r1, r2")
    program.op("jump_fwd wav_average_ready")
    program.label("wav_average_six")
    program.op("movi r3, 2")
    program.op("shl r2, r1, r3")
    program.op("movi r3, 1")
    program.op("shl r1, r1, r3")
    program.op("add r1, r1, r2")
    program.label("wav_average_ready")
    _load(program, 2, WAV_STATE["average_rate"])
    program.op("cmp r2, r1")
    program.op("jump_fwd wav_bad_format if=ne")
    program.op("movi r1, 1")
    _store(program, 1, WAV_STATE["fmt_seen"])
    program.op("jump_fwd wav_chunk_padding")

    program.label("wav_data_chunk")
    _load(program, 5, WAV_STATE["chunk_size"])
    _load(program, 1, WAV_STATE["fmt_seen"])
    _load(program, 2, WAV_STATE["data_seen"])
    program.op("movi r3, 1")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_required_chunk if=ne")
    program.op("cmp r2, r0")
    program.op("jump_fwd wav_required_chunk if=ne")
    _load(program, 1, WAV_STATE["align"])
    program.op("mov r2, r5")
    program.op("call_fwd wav_validate_alignment")
    program.op("movi r1, 1")
    _store(program, 1, WAV_STATE["data_seen"])
    _load(program, 5, WAV_STATE["chunk_size"])
    _store(program, 5, WAV_STATE["data_remaining"])
    program.op("call_fwd wav_validate_job_geometry")
    _load(program, 1, WAV_STATE["profile"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_data_passthrough if=eq")
    program.op("call_fwd wav_process_resampled")
    program.op("jump_fwd wav_data_done")
    program.label("wav_data_passthrough")
    _load(program, 1, WAV_STATE["data_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_data_done if=eq")
    _load(program, 2, WAV_STATE["align"])
    program.op("movi r3, 0")
    program.op("movi r4, 65535")
    program.op("min r4, r1, r4")
    program.op("loop_setup 1, r4")
    program.label("wav_tile_count_loop")
    program.op("cmp r1, r2")
    program.op("sub r1, r1, r2 if=uge")
    program.op("movi r5, 1")
    program.op("add r3, r3, r5 if=uge")
    program.op("loop_back 1, wav_tile_count_loop")
    program.op("cmp r1, r0")
    program.op("movi r1, 65535 if=ne")
    program.op("jump_fwd wav_tile_count_ready if=ne")
    program.op("movi r1, 31")
    program.op("add r3, r3, r1")
    program.op("movi r1, 5")
    program.op("shr r1, r3, r1")
    program.label("wav_tile_count_ready")
    program.op("loop_setup 1, r1")
    program.label("wav_data_tile_loop")
    _load(program, 1, WAV_STATE["data_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_data_tile_idle if=eq")
    _load(program, 1, WAV_STATE["data_remaining"])
    # Select min(32, remaining/alignment) without division. Stereo S24 then
    # fills the frozen 256-byte maximum output command exactly.
    program.op("movi r2, 0")
    _load(program, 3, WAV_STATE["align"])
    program.op("mov r4, r1")
    program.op("movi r5, 32")
    program.op("loop_setup 2, r5")
    program.label("wav_tile_size_loop")
    program.op("cmp r4, r3")
    program.op("sub r4, r4, r3 if=uge")
    program.op("movi r6, 1")
    program.op("add r2, r2, r6 if=uge")
    program.op("loop_back 2, wav_tile_size_loop")
    program.op("cmp r2, r0")
    program.op("jump_fwd wav_bad_chunk if=eq")
    _store(program, 2, WAV_STATE["tile_frames"])
    _store(program, 2, WAV_STATE["source_tile_frames"])
    program.op("movi r1, 0")
    program.op("loop_setup 2, r2")
    program.label("wav_read_tile_loop")
    program.op("call_fwd wav_read_frame")
    program.op("loop_back 2, wav_read_tile_loop")
    _store(program, 0, WAV_STATE["pack_input"])
    program.op("movi r1, 32")
    _store(program, 1, WAV_STATE["pack_capacity"])
    _load(program, 1, WAV_STATE["source_tile_frames"])
    _load(program, 2, WAV_STATE["align"])
    # Multiply the bounded source-frame count by the validated block alignment.
    program.op("movi r3, 0")
    program.op("loop_setup 2, r1")
    program.label("wav_tile_input_loop")
    program.op("add r3, r3, r2")
    program.op("loop_back 2, wav_tile_input_loop")
    _load(program, 4, WAV_STATE["uncommitted_input"])
    program.op("add r4, r4, r3")
    _store(program, 4, WAV_STATE["uncommitted_input"])
    _load(program, 4, WAV_STATE["data_remaining"])
    program.op("sub r4, r4, r3")
    _store(program, 4, WAV_STATE["data_remaining"])
    program.op("call_fwd wav_pack_tile")
    program.label("wav_data_tile_idle")
    program.op("loop_back 1, wav_data_tile_loop")
    # More than 65535 32-frame tiles is outside the bounded release program.
    _load(program, 1, WAV_STATE["data_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_profile_limit if=ne")
    program.label("wav_data_done")

    program.label("wav_chunk_padding")
    _load(program, 2, WAV_STATE["chunk_pad"])
    program.op("cmp r2, r0")
    program.op("jump_fwd wav_chunk_idle if=eq")
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("call_fwd wav_get_byte")
    _load(program, 1, WAV_STATE["strict"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_relaxed_pad if=eq")
    program.op("cmp r15, r0")
    program.op("jump_fwd wav_bad_pad if=ne")
    program.op("jump_fwd wav_chunk_idle")
    program.label("wav_relaxed_pad")
    program.op("cmp r15, r0")
    program.op("jump_fwd wav_chunk_idle if=eq")
    _load(program, 1, WAV_STATE["warnings"])
    program.op("movi r2, 4")
    program.op("or r1, r1, r2")
    _store(program, 1, WAV_STATE["warnings"])
    program.label("wav_chunk_idle")
    program.op("loop_back 0, wav_chunk_loop")

    _load(program, 1, WAV_STATE["riff_end"])
    program.op("cmp r11, r1")
    program.op("jump_fwd wav_profile_limit if=ne")
    program.label("wav_chunks_complete")
    _load(program, 1, WAV_STATE["fmt_seen"])
    _load(program, 2, WAV_STATE["data_seen"])
    program.op("and r1, r1, r2")
    program.op("movi r2, 1")
    program.op("cmp r1, r2")
    program.op("jump_fwd wav_required_chunk if=ne")
    _load(program, 1, WAV_STATE["strict"])
    _load(program, 2, WAV_STATE["input_length"])
    _load(program, 3, WAV_STATE["riff_end"])
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_trailing_relaxed if=ne")
    program.op("jump_fwd wav_emit_final")
    program.label("wav_trailing_relaxed")
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_bad_trailing if=ne")
    _load(program, 1, WAV_STATE["warnings"])
    program.op("movi r2, 2")
    program.op("or r1, r1, r2")
    _store(program, 1, WAV_STATE["warnings"])
    program.label("wav_emit_final")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_success if=eq")
    program.op("movi r1, 1")
    _store(program, 1, CTX_FINAL)
    program.op("call_fwd wav_emit_pending")
    program.label("wav_success")
    _load(program, 5, WAV_STATE["warnings"])
    program.op("movi r1, 16")
    program.op("shl r5, r5, r1")
    program.op("movi r4, 0")
    program.op("job_result r4, r5, 0")
    program.op("end")

    # Helpers are after all callers because V1 CALL targets are forward-only.
    program.label("wav_validate_alignment")
    # r1=alignment, r2=size. Power-of-two alignments use masks; 3/6 use
    # byte-sum modulo three plus the even check for six.
    program.op("movi r3, 1")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_alignment_valid if=eq")
    program.op("movi r3, 2")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_alignment_power if=eq")
    program.op("movi r3, 4")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_alignment_power if=eq")
    program.op("movi r3, 8")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_alignment_power if=eq")
    program.op("movi r3, 3")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_alignment_three if=eq")
    program.op("movi r3, 6")
    program.op("cmp r1, r3")
    program.op("jump_fwd wav_bad_format if=ne")
    program.op("movi r3, 1")
    program.op("and r4, r2, r3")
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_bad_chunk if=ne")
    program.label("wav_alignment_three")
    program.op("movi r3, 0xff")
    program.op("and r4, r2, r3")
    program.op("movi r3, 8")
    program.op("shr r5, r2, r3")
    program.op("movi r3, 0xff")
    program.op("and r5, r5, r3")
    program.op("add r4, r4, r5")
    program.op("movi r3, 16")
    program.op("shr r5, r2, r3")
    program.op("movi r3, 0xff")
    program.op("and r5, r5, r3")
    program.op("add r4, r4, r5")
    program.op("movi r3, 24")
    program.op("shr r5, r2, r3")
    program.op("add r4, r4, r5")
    program.op("movi r3, 340")
    program.op("loop_setup 2, r3")
    program.label("wav_modulo_three_loop")
    program.op("movi r3, 3")
    program.op("cmp r4, r3")
    program.op("sub r4, r4, r3 if=uge")
    program.op("loop_back 2, wav_modulo_three_loop")
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_bad_chunk if=ne")
    program.op("ret")
    program.label("wav_alignment_power")
    program.op("movi r3, 1")
    program.op("sub r3, r1, r3")
    program.op("and r4, r2, r3")
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_bad_chunk if=ne")
    program.label("wav_alignment_valid")
    program.op("ret")

    program.label("wav_validate_job_geometry")
    # Derive and publish source information, then enforce expected fields and
    # source-dependent output/downmix/resampling rules.
    _load(program, 1, WAV_STATE["rate"])
    _load(program, 2, WAV_STATE["channels"])
    _load(program, 3, WAV_STATE["bits"])
    program.op("movi r4, 17")
    program.op("shl r5, r2, r4")
    program.op("or r1, r1, r5")
    program.op("movi r4, 19")
    program.op("shl r5, r3, r4")
    program.op("or r1, r1, r5")
    _store(program, 1, CTX_SOURCE_INFO)
    _load(program, 4, CTX_INPUT_CONFIG)
    program.op("movi r5, 0x1ffff")
    program.op("and r5, r4, r5")
    program.op("cmp r5, r0")
    program.op("jump_fwd wav_expected_channels if=eq")
    _load(program, 6, WAV_STATE["rate"])
    program.op("cmp r5, r6")
    program.op("jump_fwd wav_source_unsupported if=ne")
    program.label("wav_expected_channels")
    program.op("movi r5, 17")
    program.op("shr r5, r4, r5")
    program.op("movi r6, 3")
    program.op("and r5, r5, r6")
    program.op("cmp r5, r0")
    program.op("jump_fwd wav_expected_bits if=eq")
    _load(program, 6, WAV_STATE["channels"])
    program.op("cmp r5, r6")
    program.op("jump_fwd wav_source_unsupported if=ne")
    program.label("wav_expected_bits")
    program.op("movi r5, 20")
    program.op("shr r5, r4, r5")
    program.op("movi r6, 0x3f")
    program.op("and r5, r5, r6")
    program.op("cmp r5, r0")
    program.op("jump_fwd wav_derive_output if=eq")
    _load(program, 6, WAV_STATE["bits"])
    program.op("cmp r5, r6")
    program.op("jump_fwd wav_source_unsupported if=ne")
    program.label("wav_derive_output")
    _load(program, 4, CTX_CONTROL)
    program.op("movi r5, 8")
    program.op("shr r5, r4, r5")
    program.op("movi r6, 3")
    program.op("and r5, r5, r6")
    _store(program, 5, WAV_STATE["output_mode"])
    _load(program, 4, CTX_OUTPUT_CONFIG)
    program.op("movi r5, 17")
    program.op("shr r5, r4, r5")
    program.op("movi r6, 3")
    program.op("and r5, r5, r6")
    program.op("cmp r5, r0")
    program.op("jump_fwd wav_output_channels_given if=ne")
    _load(program, 6, WAV_STATE["output_mode"])
    program.op("cmp r6, r0")
    program.op("jump_fwd wav_output_channels_memory if=eq")
    program.op("movi r5, 2")
    program.op("jump_fwd wav_output_channels_given")
    program.label("wav_output_channels_memory")
    _load(program, 5, WAV_STATE["channels"])
    program.label("wav_output_channels_given")
    _store(program, 5, WAV_STATE["output_channels"])
    program.op("movi r6, 19")
    program.op("shr r6, r4, r6")
    program.op("movi r7, 3")
    program.op("and r6, r6, r7")
    program.op("movi r7, 16")
    program.op("cmp r6, r0")
    program.op("jump_fwd wav_output_bits_ready if=eq")
    program.op("movi r7, 24")
    program.label("wav_output_bits_ready")
    _store(program, 7, WAV_STATE["output_bits"])
    program.op("movi r6, 0x1ffff")
    program.op("and r6, r4, r6")
    program.op("cmp r6, r0")
    program.op("jump_fwd wav_output_rate_derived if=ne")
    _load(program, 6, WAV_STATE["rate"])
    program.label("wav_output_rate_derived")
    _store(program, 6, WAV_STATE["output_rate"])
    _load(program, 4, CTX_CONTROL)
    program.op("movi r6, 1")
    program.op("movi r7, 10")
    program.op("shl r6, r6, r7")
    program.op("and r6, r4, r6")
    _store(program, 6, WAV_STATE["downmix"])
    program.op("movi r6, 1")
    program.op("movi r7, 11")
    program.op("shl r6, r6, r7")
    program.op("and r6, r4, r6")
    _store(program, 6, WAV_STATE["resample"])
    _load(program, 1, WAV_STATE["channels"])
    _load(program, 2, WAV_STATE["output_channels"])
    _load(program, 3, WAV_STATE["downmix"])
    program.op("cmp r3, r0")
    program.op("jump_fwd wav_no_downmix if=eq")
    program.op("movi r4, 2")
    program.op("cmp r1, r4")
    program.op("jump_fwd wav_output_unsupported if=ne")
    program.op("movi r4, 1")
    program.op("cmp r2, r4")
    program.op("jump_fwd wav_output_unsupported if=ne")
    program.op("jump_fwd wav_rate_check")
    program.label("wav_no_downmix")
    program.op("movi r4, 2")
    program.op("cmp r1, r4")
    program.op("jump_fwd wav_rate_check if=ne")
    program.op("movi r4, 1")
    program.op("cmp r2, r4")
    program.op("jump_fwd wav_output_unsupported if=eq")
    program.label("wav_rate_check")
    _load(program, 1, WAV_STATE["resample"])
    _load(program, 2, WAV_STATE["rate"])
    _load(program, 3, WAV_STATE["output_rate"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_rate_passthrough if=eq")
    program.op("movi r4, 48000")
    program.op("cmp r3, r4")
    program.op("jump_fwd wav_resample_profile if=eq")
    program.op("movi r4, 96000")
    program.op("cmp r3, r4")
    program.op("jump_fwd wav_output_unsupported if=ne")
    program.label("wav_resample_profile")
    # Dispatch by source once, then distinguish the already-validated 48/96
    # kHz destination. This covers the complete frozen ratio table compactly.
    profiles_by_source = {
        8000: (8, 10),
        11025: (None, 13),
        12000: (7, 9),
        16000: (6, 8),
        22050: (None, 12),
        24000: (5, 7),
        32000: (4, 6),
        44100: (3, 11),
        48000: (0, 5),
        72000: (14, 15),
        88200: (2, None),
        96000: (1, 0),
    }
    for source_rate in profiles_by_source:
        program.op(f"movi r4, {source_rate}")
        program.op("cmp r2, r4")
        program.op(f"jump_fwd wav_profile_source_{source_rate} if=eq")
    program.op("jump_fwd wav_output_unsupported")
    for source_rate, (profile_48, profile_96) in profiles_by_source.items():
        program.label(f"wav_profile_source_{source_rate}")
        if profile_48 is None:
            program.op("movi r4, 96000")
            program.op("cmp r3, r4")
            program.op("jump_fwd wav_output_unsupported if=ne")
            program.op(f"movi r4, {profile_96}")
        elif profile_96 is None:
            program.op("movi r4, 48000")
            program.op("cmp r3, r4")
            program.op("jump_fwd wav_output_unsupported if=ne")
            program.op(f"movi r4, {profile_48}")
        else:
            program.op("movi r4, 48000")
            program.op("cmp r3, r4")
            program.op(f"movi r4, {profile_48} if=eq")
            program.op(f"movi r5, {profile_96}")
            program.op("mov r4, r5 if=ne")
        program.op("jump_fwd wav_profile_record")
    program.op("nop")
    program.label("wav_profile_record")
    _store(program, 4, WAV_STATE["profile"])
    program.op("jump_fwd wav_physical_channels")
    program.label("wav_rate_passthrough")
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_output_unsupported if=ne")
    program.label("wav_physical_channels")
    _load(program, 1, WAV_STATE["output_mode"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_resampler_check if=eq")
    program.op("movi r1, 2")
    _store(program, 1, WAV_STATE["output_channels"])
    program.label("wav_resampler_check")
    _load(program, 4, WAV_STATE["profile"])
    program.op("cmp r4, r0")
    program.op("jump_fwd wav_init_resampler if=ne")
    program.op("ret")
    program.label("wav_init_resampler")
    for offset, value in (
        (0x4500, 0x80000000),
        (0x4504, 0x80000800),
        (0x4508, 0x80001000),
        (0x450C, 0),
        (0x4510, 0),
        (0x4514, 0),
        (0x4518, 0),
    ):
        program.op(f"movi r4, 0x{value:08x}")
        _store(program, 4, offset)
    _load(program, 4, WAV_STATE["output_channels"])
    _store(program, 4, 0x451C)
    program.op("ret")

    program.label("wav_process_resampled")
    # Valid zero history/look-ahead surrounds an interleaved eight-frame
    # current tile. The window is shifted and extended after each call.
    program.op("movi r1, 0")
    program.op("movi r2, 48")
    program.op("loop_setup 1, r2")
    program.label("wav_resample_zero_loop")
    program.op("st32 r1, r0, 0")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.op("loop_back 1, wav_resample_zero_loop")
    _load(program, 1, WAV_STATE["output_channels"])
    program.op("movi r2, 2")
    program.op("shl r1, r1, r2")
    program.op("movi r2, 7")
    program.op("loop_setup 1, r2")
    program.label("wav_resample_history_offset")
    program.op("add r0, r0, r1")
    program.op("loop_back 1, wav_resample_history_offset")
    program.op("mov r1, r0")
    program.op("movi r0, 0")
    program.op("movi r2, 17")
    program.op("loop_setup 1, r2")
    program.label("wav_resample_initial_loop")
    _load(program, 2, WAV_STATE["data_remaining"])
    _load(program, 3, WAV_STATE["align"])
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_resample_initial_idle if=ult")
    program.op("call_fwd wav_read_frame")
    _load(program, 2, WAV_STATE["data_remaining"])
    _load(program, 3, WAV_STATE["align"])
    program.op("sub r2, r2, r3")
    _store(program, 2, WAV_STATE["data_remaining"])
    _load(program, 4, WAV_STATE["queued_frames"])
    program.op("movi r5, 1")
    program.op("add r4, r4, r5")
    _store(program, 4, WAV_STATE["queued_frames"])
    program.label("wav_resample_initial_idle")
    program.op("loop_back 1, wav_resample_initial_loop")
    _load(program, 1, WAV_STATE["queued_frames"])
    _load(program, 2, WAV_STATE["data_remaining"])
    program.op("add r1, r1, r2")
    program.op("movi r2, 65535")
    program.op("min r1, r1, r2")
    program.op("loop_setup 1, r1")
    program.label("wav_resample_tile_loop")
    _load(program, 1, WAV_STATE["queued_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_resample_tile_idle if=eq")
    program.op("movi r2, 8")
    program.op("min r2, r1, r2")
    _store(program, 2, WAV_STATE["source_tile_frames"])
    _load(program, 2, WAV_STATE["source_tile_frames"])
    _load(program, 3, WAV_STATE["output_channels"])
    program.op("movi r4, 2")
    program.op("shl r3, r3, r4")
    program.op("movi r1, 0")
    program.op("movi r4, 7")
    program.op("loop_setup 2, r4")
    program.label("wav_resample_input_offset")
    program.op("add r1, r1, r3")
    program.op("loop_back 2, wav_resample_input_offset")
    program.op("movi r4, 0x0800")
    _store(program, 0, WAV_STATE["tile_frames"])
    program.op("loop_setup 2, r2")
    program.label("wav_resample_frame_loop")
    program.op("mov r8, r4")
    program.op("movi r2, 0x4500")
    program.op("call_fwd wav_resample_one")
    _load(program, 5, WAV_STATE["tile_frames"])
    program.op("add r5, r5, r8")
    _store(program, 5, WAV_STATE["tile_frames"])
    _load(program, 6, WAV_STATE["output_channels"])
    program.op("movi r7, 2")
    program.op("shl r6, r6, r7")
    program.op("shl r7, r8, r7")
    program.op("movi r8, 1")
    _load(program, 8, WAV_STATE["output_channels"])
    program.op("movi r2, 2")
    program.op("cmp r8, r2")
    program.op("movi r2, 1")
    program.op("shl r7, r7, r2 if=eq")
    program.label("wav_resample_output_advance")
    program.op("add r4, r4, r7")
    program.op("add r1, r1, r6")
    program.op("loop_back 2, wav_resample_frame_loop")
    program.op("movi r1, 0x0800")
    _store(program, 1, WAV_STATE["pack_input"])
    program.op("movi r1, 96")
    _store(program, 1, WAV_STATE["pack_capacity"])
    _load(program, 2, WAV_STATE["align"])
    program.op("movi r3, 0")
    _load(program, 5, WAV_STATE["source_tile_frames"])
    program.op("loop_setup 2, r5")
    program.label("wav_resample_input_account_loop")
    program.op("add r3, r3, r2")
    program.op("loop_back 2, wav_resample_input_account_loop")
    _load(program, 6, WAV_STATE["uncommitted_input"])
    program.op("add r6, r6, r3")
    _store(program, 6, WAV_STATE["uncommitted_input"])
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_resample_no_output if=eq")
    program.op("call_fwd wav_pack_tile")
    program.label("wav_resample_no_output")
    # Shift sixteen interleaved history/current/look-ahead frames left by the
    # source tile count, then append the same number of new source frames.
    _load(program, 1, WAV_STATE["output_channels"])
    program.op("movi r2, 2")
    program.op("shl r1, r1, r2")
    _load(program, 3, WAV_STATE["source_tile_frames"])
    program.op("movi r2, 0")
    program.op("loop_setup 2, r3")
    program.label("wav_resample_shift_source_loop")
    program.op("add r2, r2, r1")
    program.op("loop_back 2, wav_resample_shift_source_loop")
    program.op("mov r3, r2")
    program.op("movi r2, 0")
    program.op("movi r4, 16")
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_output_unsupported if=eq")
    program.op("movi r6, 1")
    program.op("cmp r1, r6")
    program.op("jump_fwd wav_resample_copy_count if=eq")
    program.op("movi r6, 1")
    program.op("shl r4, r4, r6")
    program.label("wav_resample_copy_count")
    program.op("loop_setup 2, r4")
    program.label("wav_resample_shift_loop")
    program.op("ld32 r5, r3, 0")
    program.op("st32 r2, r5, 0")
    program.op("movi r6, 4")
    program.op("add r2, r2, r6")
    program.op("add r3, r3, r6")
    program.op("loop_back 2, wav_resample_shift_loop")
    _load(program, 1, WAV_STATE["output_channels"])
    program.op("movi r2, 6")
    program.op("shl r1, r1, r2")
    program.label("wav_resample_append_pointer")
    _load(program, 5, WAV_STATE["source_tile_frames"])
    program.op("loop_setup 2, r5")
    program.label("wav_resample_append_loop")
    _load(program, 2, WAV_STATE["data_remaining"])
    _load(program, 3, WAV_STATE["align"])
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_resample_append_idle if=ult")
    program.op("call_fwd wav_read_frame")
    _load(program, 2, WAV_STATE["data_remaining"])
    _load(program, 3, WAV_STATE["align"])
    program.op("sub r2, r2, r3")
    _store(program, 2, WAV_STATE["data_remaining"])
    _load(program, 4, WAV_STATE["queued_frames"])
    program.op("movi r6, 1")
    program.op("add r4, r4, r6")
    _store(program, 4, WAV_STATE["queued_frames"])
    program.op("jump_fwd wav_resample_append_next")
    program.label("wav_resample_append_idle")
    program.op("st32 r1, r0, 0")
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 1")
    program.op("cmp r2, r3")
    program.op("jump_fwd wav_resample_append_advance if=eq")
    program.op("st32 r1, r0, 4")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.label("wav_resample_append_advance")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.label("wav_resample_append_next")
    program.op("loop_back 2, wav_resample_append_loop")
    _load(program, 1, WAV_STATE["queued_frames"])
    _load(program, 5, WAV_STATE["source_tile_frames"])
    program.op("sub r1, r1, r5")
    _store(program, 1, WAV_STATE["queued_frames"])
    program.label("wav_resample_tile_idle")
    program.op("loop_back 1, wav_resample_tile_loop")
    _load(program, 1, WAV_STATE["queued_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_profile_limit if=ne")
    program.op("ret")

    program.label("wav_resample_one")
    _load(program, 3, WAV_STATE["profile"])
    for profile in range(1, 16):
        label = f"wav_resample_one_{profile}"
        program.op(f"movi r6, {profile}")
        program.op("cmp r3, r6")
        program.op(f"jump_fwd {label} if=eq")
    program.op("jump_fwd wav_output_unsupported")
    for profile in range(1, 16):
        program.label(f"wav_resample_one_{profile}")
        program.op(f"resample r8, r1, r2, {profile}, 1")
        program.op("jump_fwd wav_resample_one_return")
    program.op("nop")
    program.label("wav_resample_one_return")
    program.op("ret")

    program.label("wav_read_frame")
    # r1 is the interleaved Q1.31 destination pointer and is preserved.
    _store(program, 1, WAV_STATE["frame_pointer"])
    _load(program, 2, WAV_STATE["bits"])
    _load(program, 3, WAV_STATE["channels"])
    program.op("movi r6, 8")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_left_8 if=eq")
    program.op("movi r6, 16")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_left_16 if=eq")
    program.op("movi r6, 24")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_left_24 if=eq")
    program.op("call_fwd wav_get_u32le")
    program.op("mov r4, r15")
    program.op("jump_fwd wav_read_right")
    program.label("wav_read_left_8")
    program.op("call_fwd wav_get_byte")
    program.op("movi r6, 128")
    program.op("sub r4, r15, r6")
    program.op("movi r6, 24")
    program.op("shl r4, r4, r6")
    program.op("jump_fwd wav_read_right")
    program.label("wav_read_left_16")
    program.op("call_fwd wav_get_u16le")
    program.op("movi r6, 16")
    program.op("shl r4, r15, r6")
    program.op("jump_fwd wav_read_right")
    program.label("wav_read_left_24")
    program.op("call_fwd wav_get_u24le")
    program.op("movi r6, 8")
    program.op("shl r4, r15, r6")
    program.label("wav_read_right")
    _load(program, 2, WAV_STATE["bits"])
    _load(program, 3, WAV_STATE["channels"])
    program.op("movi r6, 1")
    program.op("cmp r3, r6")
    program.op("jump_fwd wav_frame_channels_ready if=eq")
    program.op("movi r6, 8")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_right_8 if=eq")
    program.op("movi r6, 16")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_right_16 if=eq")
    program.op("movi r6, 24")
    program.op("cmp r2, r6")
    program.op("jump_fwd wav_read_right_24 if=eq")
    program.op("call_fwd wav_get_u32le")
    program.op("mov r5, r15")
    program.op("jump_fwd wav_frame_channels_ready")
    program.label("wav_read_right_8")
    program.op("call_fwd wav_get_byte")
    program.op("movi r6, 128")
    program.op("sub r5, r15, r6")
    program.op("movi r6, 24")
    program.op("shl r5, r5, r6")
    program.op("jump_fwd wav_frame_channels_ready")
    program.label("wav_read_right_16")
    program.op("call_fwd wav_get_u16le")
    program.op("movi r6, 16")
    program.op("shl r5, r15, r6")
    program.op("jump_fwd wav_frame_channels_ready")
    program.label("wav_read_right_24")
    program.op("call_fwd wav_get_u24le")
    program.op("movi r6, 8")
    program.op("shl r5, r15, r6")
    program.label("wav_frame_channels_ready")
    _load(program, 6, WAV_STATE["downmix"])
    program.op("cmp r6, r0")
    program.op("jump_fwd wav_frame_no_downmix if=eq")
    program.op("movi r6, 1")
    program.op("sar r7, r4, r6")
    program.op("sar r8, r5, r6")
    program.op("add r7, r7, r8")
    program.op("and r8, r4, r6")
    program.op("and r6, r5, r6")
    program.op("add r8, r8, r6")
    program.op("movi r6, 1")
    program.op("cmp r8, r6")
    program.op("jump_fwd wav_frame_average_ready if=ne")
    program.op("and r6, r7, r6")
    program.op("add r7, r7, r6")
    program.label("wav_frame_average_ready")
    program.op("mov r4, r7")
    program.op("mov r5, r7")
    program.op("jump_fwd wav_frame_store")
    program.label("wav_frame_no_downmix")
    _load(program, 6, WAV_STATE["channels"])
    _load(program, 7, WAV_STATE["output_channels"])
    program.op("movi r8, 1")
    program.op("cmp r6, r8")
    program.op("jump_fwd wav_frame_store if=ne")
    program.op("movi r8, 2")
    program.op("cmp r7, r8")
    program.op("mov r5, r4 if=eq")
    program.label("wav_frame_store")
    _load(program, 1, WAV_STATE["frame_pointer"])
    program.op("st32 r1, r4, 0")
    _load(program, 6, WAV_STATE["output_channels"])
    program.op("movi r7, 1")
    program.op("cmp r6, r7")
    program.op("jump_fwd wav_frame_advance if=eq")
    program.op("st32 r1, r5, 4")
    program.op("movi r7, 4")
    program.op("add r1, r1, r7")
    program.label("wav_frame_advance")
    program.op("movi r7, 4")
    program.op("add r1, r1, r7")
    _store(program, 1, WAV_STATE["frame_pointer"])
    program.op("ret")

    program.label("wav_pack_tile")
    # Keep one complete command pending for final-marker lookahead. A new tile
    # proves a full prior command is nonfinal and may be retired.
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_pack_preflush_done if=eq")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("movi r3, 256")
    program.op("cmp r2, r3")
    program.op("call_fwd wav_emit_pending if=eq")
    _load(program, 1, WAV_STATE["output_mode"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_pack_preflush_done if=ne")
    _load(program, 1, CTX_OUTPUT_CAPACITY)
    _load(program, 2, CTX_OUTPUT_CURSOR)
    program.op("sub r1, r1, r2")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("cmp r1, r2")
    program.op("call_fwd wav_emit_pending if=eq")
    program.label("wav_pack_preflush_done")
    _load(program, 1, WAV_STATE["output_mode"])
    program.op("cmp r1, r0")
    program.op("call_fwd p5_limit_pack if=eq")
    _load(program, 1, WAV_STATE["pending_input"])
    _load(program, 2, WAV_STATE["uncommitted_input"])
    program.op("add r1, r1, r2")
    _store(program, 1, WAV_STATE["pending_input"])
    _store(program, 0, WAV_STATE["uncommitted_input"])
    _load(program, 1, WAV_STATE["tile_frames"])
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 1")
    program.op("cmp r2, r3")
    program.op("shl r1, r1, r3 if=ne")
    _load(program, 3, WAV_STATE["pack_input"])
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r6, 16")
    program.op("movi r7, 8")
    program.op("cmp r4, r6")
    program.op("mov r7, r6 if=eq")
    program.op("loop_setup 2, r1")
    program.label("wav_pack_shift_loop")
    program.op("ld32 r5, r3, 0")
    program.op("sar r5, r5, r7")
    program.op("st32 r3, r5, 0")
    program.op("movi r6, 4")
    program.op("add r3, r3, r6")
    program.op("loop_back 2, wav_pack_shift_loop")
    program.op("movi r3, 0x40000000")
    program.op("st32 r0, r3, 0x4540")
    program.op("st32 r0, r2, 0x4544")
    program.op("movi r1, 0x4570")
    _store(program, 1, WAV_STATE["pending_output"])
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("loop_setup 3, r1")
    program.label("wav_pack_frame_loop")
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_pack_frame_idle if=eq")
    _load(program, 6, WAV_STATE["pending_bytes"])
    program.op("movi r7, 256")
    program.op("cmp r6, r7")
    program.op("call_fwd wav_emit_pending if=eq")
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("movi r6, 8")
    program.op("min r6, r1, r6")
    program.op("movi r7, 256")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("sub r7, r7, r2")
    program.op("movi r3, 1")
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r5, 16")
    program.op("cmp r4, r5")
    program.op("movi r5, 1")
    program.op("add r3, r3, r5 if=ne")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("cmp r4, r5")
    program.op("add r3, r3, r5 if=ne")
    program.op("shr r7, r7, r3")
    program.op("min r6, r6, r7")
    _store(program, 6, WAV_STATE["pack_capacity"])
    _load(program, 1, WAV_STATE["pack_input"])
    program.op("movi r2, 0x4540")
    program.op("movi r5, 0x4570")
    _load(program, 7, WAV_STATE["pending_bytes"])
    program.op("add r5, r5, r7")
    _load(program, 4, WAV_STATE["output_bits"])
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("call_fwd p5_pack_small")
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 2")
    program.op("shl r2, r2, r3")
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("movi r7, 0")
    program.op("loop_setup 2, r6")
    program.label("wav_pack_input_advance_loop")
    program.op("add r7, r7, r2")
    program.op("loop_back 2, wav_pack_input_advance_loop")
    program.op("add r1, r1, r7")
    _store(program, 1, WAV_STATE["pack_input"])
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r4, r7")
    program.op("movi r4, 2")
    program.op("movi r7, 4")
    program.op("mov r4, r7 if=ne")
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 1")
    program.op("cmp r2, r3")
    program.op("shl r4, r4, r3 if=ne")
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("movi r7, 0")
    program.op("loop_setup 2, r6")
    program.label("wav_pack_byte_count_loop")
    program.op("add r7, r7, r4")
    program.op("loop_back 2, wav_pack_byte_count_loop")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("add r2, r2, r7")
    _store(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 2, WAV_STATE["pending_frames"])
    program.op("add r2, r2, r6")
    _store(program, 2, WAV_STATE["pending_frames"])
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("sub r1, r1, r6")
    _store(program, 1, WAV_STATE["tile_frames"])
    program.op("movi r3, 1")
    _store(program, 3, WAV_STATE["pending"])
    program.label("wav_pack_frame_idle")
    program.op("loop_back 3, wav_pack_frame_loop")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("call_fwd wav_emit_pending if=ne")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("jump_fwd p5_capacity_fail if=ne")
    program.op("ret")

    program.label("wav_emit_pending")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_emit_return if=eq")
    _load(program, 1, WAV_STATE["pending_output"])
    _load(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 3, WAV_STATE["output_mode"])
    program.op("cmp r3, r0")
    program.op("jump_fwd wav_emit_stream if=ne")
    program.op("output_commit r8, r1, r2")
    program.op("jump_fwd wav_emit_commit")
    program.label("wav_emit_stream")
    _load(program, 3, WAV_STATE["inflight"])
    program.op("cmp r3, r0")
    program.op("call_fwd wav_retire_inflight if=ne")
    _load(program, 1, WAV_STATE["pending_output"])
    _load(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 5, CTX_FINAL)
    _store(program, 5, WAV_STATE["inflight_final"])
    program.op("output_stream r10, r1, r2")
    _load(program, 3, WAV_STATE["pending_frames"])
    _store(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["pending_input"])
    _store(program, 4, WAV_STATE["inflight_input"])
    program.op("movi r3, 1")
    _store(program, 3, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["pending"])
    _store(program, 0, WAV_STATE["pending_bytes"])
    _store(program, 0, WAV_STATE["pending_frames"])
    _store(program, 0, WAV_STATE["pending_input"])
    _load(program, 1, WAV_STATE["inflight_final"])
    program.op("cmp r1, r0")
    program.op("call_fwd wav_retire_inflight if=ne")
    program.op("jump_fwd wav_emit_capacity")
    program.label("wav_emit_commit")
    program.op("dma_wait r8")
    _load(program, 3, WAV_STATE["pending_frames"])
    _load(program, 4, WAV_STATE["pending_input"])
    _load(program, 5, WAV_STATE["output_channels"])
    _load(program, 6, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r6, r7")
    program.op("movi r6, 2")
    program.op("jump_fwd wav_emit_sample_bytes if=eq")
    program.op("movi r6, 4")
    program.label("wav_emit_sample_bytes")
    program.op("movi r7, 1")
    program.op("cmp r5, r7")
    program.op("jump_fwd wav_emit_frame_bytes if=eq")
    program.op("shl r6, r6, r7")
    program.label("wav_emit_frame_bytes")
    program.op("loop_setup 2, r3")
    program.label("wav_emit_commit_loop")
    program.op("frame_commit r4, r6")
    program.op("movi r4, 0")
    program.op("loop_back 2, wav_emit_commit_loop")
    _store(program, 0, WAV_STATE["pending"])
    _store(program, 0, WAV_STATE["pending_bytes"])
    _store(program, 0, WAV_STATE["pending_frames"])
    _store(program, 0, WAV_STATE["pending_input"])
    program.label("wav_emit_capacity")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("jump_fwd p5_capacity_fail if=ne")
    program.label("wav_emit_return")
    program.op("ret")

    program.label("wav_retire_inflight")
    _load(program, 1, WAV_STATE["inflight"])
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_retire_return if=eq")
    program.op("dma_wait r8")
    _load(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["inflight_input"])
    _load(program, 5, WAV_STATE["output_channels"])
    _load(program, 6, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r6, r7")
    program.op("movi r6, 2")
    program.op("jump_fwd wav_retire_sample_bytes if=eq")
    program.op("movi r6, 4")
    program.label("wav_retire_sample_bytes")
    program.op("movi r7, 1")
    program.op("cmp r5, r7")
    program.op("jump_fwd wav_retire_frame_bytes if=eq")
    program.op("shl r6, r6, r7")
    program.label("wav_retire_frame_bytes")
    program.op("loop_setup 2, r3")
    program.label("wav_retire_commit_loop")
    program.op("frame_commit r4, r6")
    program.op("movi r4, 0")
    program.op("loop_back 2, wav_retire_commit_loop")
    _store(program, 0, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["inflight_frames"])
    _store(program, 0, WAV_STATE["inflight_input"])
    _store(program, 0, WAV_STATE["inflight_final"])
    program.label("wav_retire_return")
    program.op("ret")

    program.label("wav_skip_bytes")
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_skip_return if=eq")
    program.op("movi r2, 65535")
    program.op("min r2, r1, r2")
    program.op("loop_setup 2, r2")
    program.label("wav_skip_loop")
    program.op("call_fwd wav_get_byte")
    program.op("loop_back 2, wav_skip_loop")
    program.op("sub r1, r1, r2")
    program.op("cmp r1, r0")
    program.op("jump_fwd wav_profile_limit if=ne")
    program.label("wav_skip_return")
    program.op("ret")

    program.label("wav_get_u32le")
    program.op("call_fwd wav_get_byte")
    program.op("mov r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 8")
    program.op("shl r15, r15, r7")
    program.op("or r6, r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 16")
    program.op("shl r15, r15, r7")
    program.op("or r6, r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 24")
    program.op("shl r15, r15, r7")
    program.op("or r15, r6, r15")
    program.op("ret")
    program.label("wav_get_u24le")
    program.op("call_fwd wav_get_byte")
    program.op("mov r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 8")
    program.op("shl r15, r15, r7")
    program.op("or r6, r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 16")
    program.op("shl r15, r15, r7")
    program.op("or r15, r6, r15")
    program.op("ret")
    program.label("wav_get_u16le")
    program.op("call_fwd wav_get_byte")
    program.op("mov r6, r15")
    program.op("call_fwd wav_get_byte")
    program.op("movi r7, 8")
    program.op("shl r15, r15, r7")
    program.op("or r15, r6, r15")
    program.op("ret")
    program.label("wav_get_byte")
    program.op("cmp r13, r0")
    program.op("jump_fwd wav_have_word if=ne")
    program.op("cmp r14, r0")
    program.op("jump_fwd wav_have_batch if=ne")
    _load(program, 1, WAV_STATE["input_length"])
    program.op("cmp r11, r1")
    program.op("jump_fwd wav_truncated if=uge")
    program.op(f"movi r2, 0x{READER_STAGING:04x}")
    program.op("movi r3, 128")
    program.op("input_refill r14, r2, r3")
    program.op("dma_wait r8")
    program.op("cmp r14, r0")
    program.op("jump_fwd wav_truncated if=eq")
    program.label("wav_have_batch")
    program.op("fifo_pop r12")
    program.op("movi r8, 7")
    program.op("and r13, r13, r8")
    program.op("sub r14, r14, r13")
    program.label("wav_have_word")
    program.op("movi r8, 0xff")
    program.op("and r15, r12, r8")
    program.op("movi r8, 8")
    program.op("shr r12, r12, r8")
    program.op("movi r8, 1")
    program.op("sub r13, r13, r8")
    program.op("add r11, r11, r8")
    _store(program, 11, CTX_PARSED)
    program.op("ret")

    # Parser branches set the first failing input byte before selecting a tuple.
    for label, code, reason in (
        ("wav_bad_extent", 4, 0x0011),
        ("wav_bad_chunk", 4, 0x0013),
        ("wav_required_chunk", 4, 0x0012),
        ("wav_bad_format", 4, 0x0014),
        ("wav_unsupported_encoding", 3, 0x0003),
        ("wav_profile_limit", 3, 0x0004),
        ("wav_bad_pad", 4, 0x0019),
        ("wav_bad_trailing", 4, 0x001B),
        ("wav_truncated", 5, 0x0020),
    ):
        program.label(label)
        _fail(program, code, reason)

    program.label("wav_source_unsupported")
    _fail(program, 3, 0x0002)
    program.label("wav_output_unsupported")
    _fail(program, 3, 0x0050, "wav_fail_output")
    program.label("wav_fail")
    _load(program, 1, WAV_STATE["warnings"])
    program.op("movi r2, 16")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    _load(program, 1, WAV_STATE["format"])
    program.op("movi r2, 24")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    program.op("job_result r4, r5, 4")
    program.op("end")
    program.label("wav_fail_output")
    _load(program, 1, WAV_STATE["warnings"])
    program.op("movi r2, 16")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    _load(program, 1, WAV_STATE["format"])
    program.op("movi r2, 24")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    program.op("job_result r4, r5, 0")
    program.label("wav_terminal")
    program.op("end")


def _emit_flac(program: Program) -> None:
    """Emit the bounded native-FLAC metadata and frame driver."""

    program.label("flac")
    program.op("movi r0, 0")
    program.op("movi r1, 0x4400")
    program.op("movi r2, 208")
    program.op("loop_setup 0, r2")
    program.label("flac_state_clear_loop")
    program.op("st32 r1, r0, 0")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.op("loop_back 0, flac_state_clear_loop")
    program.op("movi r1, 2")
    _store(program, 1, WAV_STATE["format"])
    program.op("movi r1, 1")
    _store(program, 1, WAV_STATE["warnings"])
    _load(program, 1, CTX_FLAGS)
    program.op("movi r2, 1")
    program.op("and r1, r1, r2")
    _store(program, 1, WAV_STATE["strict"])
    _load(program, 1, CTX_INPUT_LENGTH)
    _store(program, 1, WAV_STATE["input_length"])
    program.op("movi r1, 4")
    _store(program, 1, FLAC_STATE["parsing_stage"])
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("movi r1, 32")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 0x664c6143")
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_bad_signature if=ne")

    program.op("movi r1, 128")
    program.op("loop_setup 0, r1")
    program.label("flac_metadata_loop")
    _load(program, 1, FLAC_STATE["metadata_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_metadata_idle if=ne")
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("mov r4, r15")
    program.op("movi r1, 7")
    program.op("shr r2, r4, r1")
    _store(program, 2, FLAC_STATE["metadata_done"])
    program.op("movi r1, 0x7f")
    program.op("and r4, r4, r1")
    _store(program, 4, FLAC_STATE["kind"])
    program.op("movi r1, 24")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["frame_remaining"])
    _load(program, 1, FLAC_STATE["metadata_count"])
    program.op("movi r2, 1")
    program.op("add r1, r1, r2")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_metadata_first if=eq")
    _load(program, 4, FLAC_STATE["kind"])
    program.op("cmp r4, r0")
    program.op("jump_fwd flac_bad_metadata if=eq")
    program.op("movi r2, 127")
    program.op("cmp r4, r2")
    program.op("jump_fwd flac_bad_metadata if=eq")
    program.op("movi r2, 7")
    program.op("cmp r4, r2")
    program.op("jump_fwd flac_unsupported_encoding if=uge")
    program.op("jump_fwd flac_metadata_account")
    program.label("flac_metadata_first")
    _load(program, 4, FLAC_STATE["kind"])
    program.op("cmp r4, r0")
    program.op("jump_fwd flac_bad_streaminfo if=ne")
    program.op("movi r2, 5")
    _store(program, 2, CTX_DIAGNOSTIC)
    _load(program, 4, FLAC_STATE["frame_remaining"])
    program.op("movi r2, 34")
    program.op("cmp r4, r2")
    program.op("jump_fwd flac_bad_streaminfo if=ne")
    program.label("flac_metadata_account")
    _store(program, 1, FLAC_STATE["metadata_count"])
    _load(program, 4, CTX_DIAGNOSTIC)
    program.op("movi r2, 1")
    program.op("add r4, r4, r2")
    _store(program, 4, CTX_DIAGNOSTIC)
    _load(program, 2, FLAC_STATE["metadata_bytes"])
    _load(program, 3, FLAC_STATE["frame_remaining"])
    program.op("add r2, r2, r3")
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    program.op("movi r3, 0x00100001")
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_profile_limit4 if=uge")
    _store(program, 2, FLAC_STATE["metadata_bytes"])
    _load(program, 4, CTX_DIAGNOSTIC)
    program.op("movi r2, 1")
    program.op("sub r4, r4, r2")
    _store(program, 4, CTX_DIAGNOSTIC)
    program.op("movi r2, 1")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_streaminfo if=eq")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("call_fwd flac_skip_bytes")
    program.op("jump_fwd flac_metadata_idle")

    program.label("flac_streaminfo")
    for width, destination in (
        (16, "min_block"),
        (16, "max_block"),
        (24, "min_frame"),
        (24, "max_frame"),
        (20, "rate"),
    ):
        program.op(f"movi r1, {width}")
        program.op("call_fwd flac_get_bits")
        _store(
            program,
            15,
            WAV_STATE[destination] if destination == "rate" else FLAC_STATE[destination],
        )
    program.op("movi r1, 3")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    _store(program, 15, WAV_STATE["channels"])
    program.op("movi r1, 5")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    _store(program, 15, WAV_STATE["bits"])
    program.op("movi r1, 4")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["total_hi"])
    program.op("movi r1, 32")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["total_lo"])
    for _ in range(4):
        program.op("movi r1, 32")
        program.op("call_fwd flac_get_bits")
    program.op("call_fwd flac_validate_streaminfo")
    program.label("flac_metadata_idle")
    program.op("loop_back 0, flac_metadata_loop")
    _load(program, 1, FLAC_STATE["metadata_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_profile_limit4 if=eq")
    program.op("call_fwd wav_validate_job_geometry")

    _load(program, 1, WAV_STATE["input_length"])
    program.op("movi r2, 65535")
    program.op("min r1, r1, r2")
    program.op("loop_setup 0, r1")
    program.label("flac_frame_loop")
    program.op("call_fwd flac_terminal_check")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_frame_idle if=ne")
    _load(program, 1, WAV_STATE["input_length"])
    program.op("cmp r11, r1")
    program.op("jump_fwd flac_sample_count if=uge")
    _store(program, 11, FLAC_STATE["frame_start"])
    _store(program, 11, CTX_DIAGNOSTIC)
    program.op("movi r6, 0")
    program.op("movi r7, 0")
    program.op("movi r8, 3")
    program.op("movi r1, 14")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 0x3ffe")
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_bad_frame if=ne")
    program.op("movi r1, 1")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_bad_frame if=ne")
    program.op("movi r1, 1")
    program.op("call_fwd flac_get_bits")
    program.op("mov r4, r15")
    _load(program, 1, FLAC_STATE["strategy_seen"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_strategy_compare if=ne")
    _store(program, 4, FLAC_STATE["strategy"])
    program.op("movi r1, 1")
    _store(program, 1, FLAC_STATE["strategy_seen"])
    program.op("jump_fwd flac_strategy_ready")
    program.label("flac_strategy_compare")
    _load(program, 1, FLAC_STATE["strategy"])
    program.op("cmp r4, r1")
    program.op("jump_fwd flac_geometry_change if=ne")
    program.label("flac_strategy_ready")
    for width, destination in (
        (4, "kind"),
        (4, "frame_remaining"),
        (4, "assignment"),
        (3, "coded_bits"),
    ):
        program.op(f"movi r1, {width}")
        program.op("call_fwd flac_get_bits")
        _store(program, 15, FLAC_STATE[destination])
    program.op("movi r1, 1")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_bad_frame if=ne")
    program.op("call_fwd flac_get_utf8")
    _store(program, 15, FLAC_STATE["parsed_number"])
    _load(program, 1, FLAC_STATE["kind"])
    program.op("call_fwd flac_decode_block_size")
    _store(program, 15, FLAC_STATE["block_size"])
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("call_fwd flac_decode_rate")
    _load(program, 1, WAV_STATE["rate"])
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_geometry_change if=ne")
    _load(program, 1, FLAC_STATE["coded_bits"])
    program.op("call_fwd flac_decode_precision")
    _load(program, 1, WAV_STATE["bits"])
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_geometry_change if=ne")
    program.op("call_fwd flac_validate_frame_header")
    program.op("movi r8, 2")
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r7")
    program.op("jump_fwd flac_header_crc if=ne")
    program.op("movi r1, 0")
    _store(program, 1, FLAC_STATE["channel"])
    _store(program, 1, FLAC_STATE["output_base"])
    program.op("call_fwd flac_decode_subframe")
    _load(program, 1, WAV_STATE["channels"])
    program.op("movi r2, 2")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_subframes_done if=ne")
    program.op("movi r1, 1")
    _store(program, 1, FLAC_STATE["channel"])
    _load(program, 1, FLAC_STATE["block_size"])
    program.op("movi r2, 2")
    program.op("shl r1, r1, r2")
    _store(program, 1, FLAC_STATE["output_base"])
    program.op("call_fwd flac_decode_subframe")
    program.label("flac_subframes_done")
    program.op("call_fwd flac_align_zero")
    program.op("movi r8, 0")
    program.op("movi r1, 16")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r6")
    program.op("jump_fwd flac_frame_crc if=ne")
    _load(program, 1, FLAC_STATE["frame_start"])
    program.op("sub r2, r11, r1")
    _store(program, 2, FLAC_STATE["frame_input"])
    # Reject advertised or observed frame sizes before decoded samples are
    # released to the conversion and output path.
    program.op("movi r3, 65537")
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_profile_limit4 if=uge")
    _load(program, 3, FLAC_STATE["min_frame"])
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_frame_max_check if=eq")
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    program.label("flac_frame_max_check")
    _load(program, 3, FLAC_STATE["max_frame"])
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_frame_profile_ready if=eq")
    program.op("cmp r3, r2")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    program.label("flac_frame_profile_ready")
    _store(program, 0, WAV_STATE["resample_final"])
    _load(program, 1, FLAC_STATE["total_lo"])
    _load(program, 2, FLAC_STATE["total_hi"])
    program.op("or r3, r1, r2")
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_final_unknown if=eq")
    _load(program, 3, FLAC_STATE["decoded_lo"])
    _load(program, 4, FLAC_STATE["block_size"])
    program.op("add r5, r3, r4")
    program.op("cmp r5, r3")
    program.op("movi r6, 1")
    _load(program, 7, FLAC_STATE["decoded_hi"])
    program.op("add r7, r7, r6 if=ult")
    program.op("cmp r2, r7")
    program.op("jump_fwd flac_sample_count if=ult")
    program.op("jump_fwd flac_final_ready if=ne")
    program.op("cmp r1, r5")
    program.op("jump_fwd flac_sample_count if=ult")
    program.op("jump_fwd flac_final_ready if=ne")
    _store(program, 6, WAV_STATE["resample_final"])
    program.op("jump_fwd flac_final_ready")
    program.label("flac_final_unknown")
    _load(program, 1, WAV_STATE["input_length"])
    program.op("cmp r11, r1")
    program.op("movi r2, 1")
    program.op(f"st32 r0, r2, 0x{WAV_STATE['resample_final']:04x} if=eq")
    program.label("flac_final_ready")
    _load(program, 1, WAV_STATE["resample_final"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_block_profile_ready if=ne")
    _load(program, 1, FLAC_STATE["block_size"])
    _load(program, 2, FLAC_STATE["min_block"])
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    program.label("flac_block_profile_ready")
    program.op("call_fwd flac_process_frame")
    _load(program, 1, FLAC_STATE["decoded_lo"])
    _load(program, 2, FLAC_STATE["block_size"])
    program.op("add r3, r1, r2")
    program.op("cmp r3, r1")
    program.op("movi r4, 1")
    _load(program, 5, FLAC_STATE["decoded_hi"])
    program.op("add r5, r5, r4 if=ult")
    _store(program, 3, FLAC_STATE["decoded_lo"])
    _store(program, 5, FLAC_STATE["decoded_hi"])
    _load(program, 1, FLAC_STATE["expected_number"])
    program.op("movi r2, 1")
    program.op("add r1, r1, r2")
    _store(program, 1, FLAC_STATE["expected_number"])
    program.label("flac_frame_idle")
    program.op("loop_back 0, flac_frame_loop")
    program.op("call_fwd flac_terminal_check")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_profile_limit4 if=eq")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_success if=eq")
    program.op("movi r1, 1")
    _store(program, 1, CTX_FINAL)
    program.op("call_fwd flac_emit_pending")
    program.label("flac_success")
    program.op("movi r4, 0")
    program.op("movi r5, 0x02010000")
    program.op("job_result r4, r5, 0")
    program.op("end")


def _emit_flac_helpers(program: Program) -> None:
    """Emit forward-only FLAC parser, reconstruction, and reader helpers."""

    program.label("flac_validate_streaminfo")
    program.op("movi r6, 8")
    _store(program, 6, CTX_DIAGNOSTIC)
    _load(program, 1, FLAC_STATE["min_block"])
    program.op("movi r2, 16")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    program.op("movi r6, 10")
    _store(program, 6, CTX_DIAGNOSTIC)
    _load(program, 2, FLAC_STATE["max_block"])
    program.op("cmp r2, r1")
    program.op("jump_fwd flac_bad_streaminfo if=ult")
    _load(program, 3, WAV_STATE["channels"])
    program.op("movi r4, 1")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_streaminfo_mono if=eq")
    program.op("movi r4, 2")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_streaminfo_frame_hints if=ne")
    program.op("movi r4, 2049")
    program.op("jump_fwd flac_streaminfo_block_limit")
    program.label("flac_streaminfo_mono")
    program.op("movi r4, 4097")
    program.label("flac_streaminfo_block_limit")
    program.op("cmp r2, r4")
    program.op("jump_fwd flac_profile_limit4 if=uge")
    program.label("flac_streaminfo_frame_hints")
    program.op("movi r6, 15")
    _store(program, 6, CTX_DIAGNOSTIC)
    _load(program, 1, FLAC_STATE["min_frame"])
    _load(program, 2, FLAC_STATE["max_frame"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_streaminfo_max_frame if=eq")
    program.op("cmp r2, r0")
    program.op("jump_fwd flac_streaminfo_max_frame if=eq")
    program.op("cmp r2, r1")
    program.op("jump_fwd flac_bad_streaminfo if=ult")
    program.label("flac_streaminfo_max_frame")
    program.op("movi r3, 65537")
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_profile_limit4 if=uge")
    program.op("movi r6, 18")
    _store(program, 6, CTX_DIAGNOSTIC)
    _load(program, 3, WAV_STATE["channels"])
    program.op("movi r4, 1")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_streaminfo_geometry_channels if=eq")
    program.op("movi r4, 2")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_unsupported_geometry if=ne")
    program.label("flac_streaminfo_geometry_channels")
    _load(program, 1, WAV_STATE["rate"])
    program.op("movi r2, 8000")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_unsupported_geometry if=ult")
    program.op("movi r2, 96001")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_unsupported_geometry if=uge")
    _load(program, 1, WAV_STATE["bits"])
    program.op("movi r2, 16")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_streaminfo_valid if=eq")
    program.op("movi r2, 24")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_unsupported_geometry if=ne")
    program.label("flac_streaminfo_valid")
    program.op("ret")

    program.label("flac_terminal_check")
    _load(program, 1, FLAC_STATE["total_lo"])
    _load(program, 2, FLAC_STATE["total_hi"])
    program.op("or r3, r1, r2")
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_terminal_unknown if=eq")
    _load(program, 3, FLAC_STATE["decoded_hi"])
    program.op("cmp r3, r2")
    program.op("jump_fwd flac_terminal_not_done if=ne")
    _load(program, 3, FLAC_STATE["decoded_lo"])
    program.op("cmp r3, r1")
    program.op("jump_fwd flac_terminal_not_done if=ne")
    program.op("movi r15, 1")
    program.op("ret")
    program.label("flac_terminal_unknown")
    _load(program, 1, WAV_STATE["input_length"])
    program.op("cmp r11, r1")
    program.op("movi r15, 1 if=eq")
    program.op("movi r15, 0 if=ne")
    program.op("ret")
    program.label("flac_terminal_not_done")
    program.op("movi r15, 0")
    program.op("ret")

    program.label("flac_decode_block_size")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_bad_frame if=eq")
    program.op("movi r2, 1")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_block_192 if=eq")
    program.op("movi r2, 6")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_block_byte if=eq")
    program.op("movi r2, 7")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_block_word if=eq")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_block_fixed if=ult")
    program.op("movi r2, 8")
    program.op("sub r2, r1, r2")
    program.op("movi r15, 256")
    program.op("shl r15, r15, r2")
    program.op("ret")
    program.label("flac_block_fixed")
    program.op("movi r2, 2")
    program.op("sub r2, r1, r2")
    program.op("movi r15, 576")
    program.op("shl r15, r15, r2")
    program.op("ret")
    program.label("flac_block_192")
    program.op("movi r15, 192")
    program.op("ret")
    program.label("flac_block_byte")
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    program.op("ret")
    program.label("flac_block_word")
    program.op("movi r1, 16")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    program.op("ret")

    program.label("flac_decode_rate")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_rate_streaminfo if=eq")
    for code, rate in (
        (1, 88200),
        (2, 176400),
        (3, 192000),
        (4, 8000),
        (5, 16000),
        (6, 22050),
        (7, 24000),
        (8, 32000),
        (9, 44100),
        (10, 48000),
        (11, 96000),
    ):
        program.op(f"movi r2, {code}")
        program.op("cmp r1, r2")
        program.op(f"jump_fwd flac_rate_{rate} if=eq")
    program.op("movi r2, 12")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_rate_khz if=eq")
    program.op("movi r2, 13")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_rate_hz if=eq")
    program.op("movi r2, 14")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_rate_ten_hz if=eq")
    program.op("jump_fwd flac_bad_frame")
    program.label("flac_rate_streaminfo")
    _load(program, 15, WAV_STATE["rate"])
    program.op("ret")
    for _, rate in (
        (1, 88200),
        (2, 176400),
        (3, 192000),
        (4, 8000),
        (5, 16000),
        (6, 22050),
        (7, 24000),
        (8, 32000),
        (9, 44100),
        (10, 48000),
        (11, 96000),
    ):
        program.label(f"flac_rate_{rate}")
        program.op(f"movi r15, {rate}")
        program.op("ret")
    program.label("flac_rate_khz")
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 10")
    program.op("shl r2, r15, r1")
    program.op("movi r1, 3")
    program.op("shl r3, r15, r1")
    program.op("movi r1, 4")
    program.op("shl r4, r15, r1")
    program.op("sub r15, r2, r3")
    program.op("sub r15, r15, r4")
    program.op("ret")
    program.label("flac_rate_hz")
    program.op("movi r1, 16")
    program.op("call_fwd flac_get_bits")
    program.op("ret")
    program.label("flac_rate_ten_hz")
    program.op("movi r1, 16")
    program.op("call_fwd flac_get_bits")
    program.op("mov r2, r15")
    program.op("movi r1, 3")
    program.op("shl r15, r15, r1")
    program.op("movi r1, 1")
    program.op("shl r2, r2, r1")
    program.op("add r15, r15, r2")
    program.op("ret")

    program.label("flac_decode_precision")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_precision_streaminfo if=eq")
    for code, bits in ((1, 8), (2, 12), (4, 16), (5, 20), (6, 24), (7, 32)):
        program.op(f"movi r2, {code}")
        program.op("cmp r1, r2")
        program.op(f"jump_fwd flac_precision_{bits} if=eq")
    program.op("jump_fwd flac_bad_frame")
    program.label("flac_precision_streaminfo")
    _load(program, 15, WAV_STATE["bits"])
    program.op("ret")
    for _, bits in ((1, 8), (2, 12), (4, 16), (5, 20), (6, 24), (7, 32)):
        program.label(f"flac_precision_{bits}")
        program.op(f"movi r15, {bits}")
        program.op("ret")

    program.label("flac_validate_frame_header")
    _load(program, 1, FLAC_STATE["assignment"])
    program.op("movi r2, 11")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_bad_frame if=uge")
    program.op("movi r2, 8")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_assignment_two if=uge")
    program.op("movi r2, 1")
    program.op("add r3, r1, r2")
    program.op("jump_fwd flac_assignment_ready")
    program.label("flac_assignment_two")
    program.op("movi r3, 2")
    program.label("flac_assignment_ready")
    _load(program, 2, WAV_STATE["channels"])
    program.op("cmp r3, r2")
    program.op("jump_fwd flac_geometry_change if=ne")
    _load(program, 1, FLAC_STATE["block_size"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_bad_frame if=eq")
    _load(program, 2, FLAC_STATE["max_block"])
    program.op("cmp r2, r1")
    program.op("jump_fwd flac_profile_limit4 if=ult")
    _load(program, 1, FLAC_STATE["strategy"])
    _load(program, 2, FLAC_STATE["parsed_number"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_number_variable if=ne")
    _load(program, 3, FLAC_STATE["expected_number"])
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_bad_number if=ne")
    program.op("ret")
    program.label("flac_number_variable")
    _load(program, 3, FLAC_STATE["decoded_lo"])
    program.op("cmp r2, r3")
    program.op("jump_fwd flac_bad_number if=ne")
    program.op("ret")

    program.label("flac_get_utf8")
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("mov r2, r15")
    program.op("movi r1, 128")
    program.op("cmp r2, r1")
    program.op("jump_fwd flac_utf8_ascii if=ult")
    program.op("movi r3, 0")
    program.op("movi r1, 0")
    program.op("movi r4, 0x80")
    program.op("movi r5, 7")
    program.op("loop_setup 1, r5")
    program.label("flac_utf8_prefix_loop")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_utf8_prefix_skip if=ne")
    program.op("and r5, r2, r4")
    program.op("cmp r5, r0")
    program.op("jump_fwd flac_utf8_prefix_zero if=eq")
    program.op("movi r5, 1")
    program.op("add r3, r3, r5")
    program.op("jump_fwd flac_utf8_prefix_skip")
    program.label("flac_utf8_prefix_zero")
    program.op("movi r1, 1")
    program.label("flac_utf8_prefix_skip")
    program.op("movi r5, 1")
    program.op("shr r4, r4, r5")
    program.op("loop_back 1, flac_utf8_prefix_loop")
    program.op("movi r4, 2")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_bad_frame if=ult")
    program.op("movi r4, 5")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_profile_limit4 if=uge")
    program.op("movi r4, 7")
    program.op("sub r4, r4, r3")
    program.op("movi r5, 1")
    program.op("shl r5, r5, r4")
    program.op("movi r4, 1")
    program.op("sub r5, r5, r4")
    program.op("and r2, r2, r5")
    _store(program, 2, FLAC_STATE["parsed_number"])
    program.op("sub r3, r3, r4")
    program.op("loop_setup 1, r3")
    program.label("flac_utf8_body_loop")
    program.op("movi r1, 8")
    program.op("call_fwd flac_get_bits")
    program.op("movi r4, 0xc0")
    program.op("and r5, r15, r4")
    program.op("movi r4, 0x80")
    program.op("cmp r5, r4")
    program.op("jump_fwd flac_bad_frame if=ne")
    program.op("movi r4, 6")
    _load(program, 2, FLAC_STATE["parsed_number"])
    program.op("shl r2, r2, r4")
    program.op("movi r4, 0x3f")
    program.op("and r15, r15, r4")
    program.op("or r2, r2, r15")
    _store(program, 2, FLAC_STATE["parsed_number"])
    program.op("loop_back 1, flac_utf8_body_loop")
    _load(program, 15, FLAC_STATE["parsed_number"])
    program.op("ret")
    program.label("flac_utf8_ascii")
    program.op("mov r15, r2")
    program.op("ret")

    program.label("flac_skip_bytes")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_skip_return if=eq")
    program.op("mov r2, r11")
    program.op("add r2, r2, r1")
    program.op("cmp r2, r11")
    program.op("jump_fwd flac_truncated if=ult")
    _load(program, 3, WAV_STATE["input_length"])
    program.op("cmp r3, r2")
    program.op("jump_fwd flac_truncated if=ult")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    # A 1 MiB aggregate can straddle the batch already holding its header.
    program.op("movi r1, 8193")
    program.op("loop_setup 2, r1")
    program.label("flac_skip_batch_loop")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_skip_batch_idle if=eq")
    program.op("cmp r13, r0")
    program.op("jump_fwd flac_skip_have_batch if=ne")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_skip_have_batch if=ne")
    program.op(f"movi r3, 0x{READER_STAGING:04x}")
    program.op("movi r4, 128")
    program.op("input_refill r14, r3, r4")
    program.op("dma_wait r5")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_truncated if=eq")
    program.label("flac_skip_have_batch")
    # Full refills are discarded a word at a time with cursor accounting once
    # per batch. The slow path below preserves a partial word at the metadata
    # boundary so the following frame/header bytes remain available.
    program.op("cmp r13, r0")
    program.op("jump_fwd flac_skip_partial_batch if=ne")
    program.op("movi r2, 128")
    program.op("cmp r14, r2")
    program.op("jump_fwd flac_skip_partial_batch if=ne")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_skip_partial_batch if=ult")
    program.op("movi r3, 32")
    program.op("loop_setup 3, r3")
    program.label("flac_skip_full_word_loop")
    program.op("fifo_pop r12")
    program.op("loop_back 3, flac_skip_full_word_loop")
    program.op("sub r1, r1, r2")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    program.op("add r11, r11, r2")
    _store(program, 11, CTX_PARSED)
    program.op("movi r13, 0")
    program.op("movi r14, 0")
    program.op("jump_fwd flac_skip_batch_idle")
    program.label("flac_skip_partial_batch")
    program.op("movi r3, 64")
    program.op("loop_setup 3, r3")
    program.label("flac_skip_word_loop")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_skip_word_idle if=eq")
    program.op("cmp r13, r0")
    program.op("jump_fwd flac_skip_have_word if=ne")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_skip_word_idle if=eq")
    program.op("fifo_pop r12")
    program.op("movi r2, 7")
    program.op("and r13, r13, r2")
    program.op("sub r14, r14, r13")
    program.label("flac_skip_have_word")
    program.op("mov r2, r13")
    program.op("min r2, r1, r2")
    program.op("movi r3, 3")
    program.op("shl r3, r2, r3")
    program.op("shr r12, r12, r3")
    program.op("sub r13, r13, r2")
    program.op("sub r1, r1, r2")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    program.op("add r11, r11, r2")
    _store(program, 11, CTX_PARSED)
    program.label("flac_skip_word_idle")
    program.op("loop_back 3, flac_skip_word_loop")
    program.label("flac_skip_batch_idle")
    program.op("loop_back 2, flac_skip_batch_loop")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_truncated if=ne")
    program.label("flac_skip_return")
    program.op("ret")

    program.label("flac_decode_subframe")
    program.op("movi r1, 5")
    _store(program, 1, FLAC_STATE["parsing_stage"])
    program.op("movi r1, 1")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_bad_entropy if=ne")
    program.op("movi r1, 6")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["kind"])
    program.op("movi r1, 1")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_no_wasted if=eq")
    program.op("call_fwd flac_get_unary")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    program.op("jump_fwd flac_wasted_ready")
    program.label("flac_no_wasted")
    program.op("movi r15, 0")
    program.label("flac_wasted_ready")
    _load(program, 1, FLAC_STATE["channel"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_store_wasted0 if=eq")
    _store(program, 15, FLAC_STATE["wasted1"])
    program.op("jump_fwd flac_coded_precision")
    program.label("flac_store_wasted0")
    _store(program, 15, FLAC_STATE["wasted0"])
    program.label("flac_coded_precision")
    _load(program, 2, WAV_STATE["bits"])
    _load(program, 3, FLAC_STATE["assignment"])
    program.op("movi r4, 8")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_side_right if=eq")
    program.op("movi r4, 9")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_side_left if=eq")
    program.op("movi r4, 10")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_side_right if=eq")
    program.op("jump_fwd flac_side_right")
    program.label("flac_side_left")
    program.op("cmp r1, r0")
    program.op("movi r4, 1")
    program.op("add r2, r2, r4 if=eq")
    program.label("flac_side_right")
    program.op("movi r4, 8")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_side_precision_mid if=ne")
    program.op("movi r4, 1")
    program.op("cmp r1, r4")
    program.op("add r2, r2, r4 if=eq")
    program.op("jump_fwd flac_side_precision_ready")
    program.label("flac_side_precision_mid")
    program.op("movi r4, 10")
    program.op("cmp r3, r4")
    program.op("jump_fwd flac_side_precision_ready if=ne")
    program.op("movi r4, 1")
    program.op("cmp r1, r4")
    program.op("add r2, r2, r4 if=eq")
    program.label("flac_side_precision_ready")
    program.op("sub r2, r2, r15")
    program.op("cmp r2, r0")
    program.op("jump_fwd flac_bad_wasted if=eq")
    _store(program, 2, FLAC_STATE["coded_bits"])
    _load(program, 1, FLAC_STATE["kind"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_subframe_constant if=eq")
    program.op("movi r2, 1")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_subframe_verbatim if=eq")
    program.op("movi r2, 8")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_subframe_fixed if=uge")
    program.op("jump_fwd flac_bad_entropy")

    program.label("flac_subframe_constant")
    _load(program, 1, FLAC_STATE["coded_bits"])
    program.op("call_fwd flac_get_signed")
    _load(program, 1, FLAC_STATE["output_base"])
    _load(program, 2, FLAC_STATE["block_size"])
    program.op("loop_setup 1, r2")
    program.label("flac_constant_loop")
    program.op("st32 r1, r15, 0")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    program.op("loop_back 1, flac_constant_loop")
    program.op("ret")

    program.label("flac_subframe_verbatim")
    _load(program, 1, FLAC_STATE["output_base"])
    _store(program, 1, FLAC_STATE["reconstruct_output"])
    _load(program, 2, FLAC_STATE["block_size"])
    program.op("loop_setup 1, r2")
    program.label("flac_verbatim_loop")
    _load(program, 1, FLAC_STATE["coded_bits"])
    program.op("call_fwd flac_get_signed")
    _load(program, 1, FLAC_STATE["reconstruct_output"])
    program.op("st32 r1, r15, 0")
    program.op("movi r3, 4")
    program.op("add r1, r1, r3")
    _store(program, 1, FLAC_STATE["reconstruct_output"])
    program.op("loop_back 1, flac_verbatim_loop")
    program.op("ret")

    program.label("flac_subframe_fixed")
    program.op("movi r2, 13")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_subframe_lpc if=uge")
    program.op("movi r2, 8")
    program.op("sub r1, r1, r2")
    program.op("jump_fwd flac_predictor_ready")
    program.label("flac_subframe_lpc")
    program.op("movi r2, 32")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_bad_entropy if=ult")
    program.op("movi r2, 64")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_bad_entropy if=uge")
    program.op("movi r2, 31")
    program.op("sub r1, r1, r2")
    program.label("flac_predictor_ready")
    _store(program, 1, FLAC_STATE["order"])
    _load(program, 2, FLAC_STATE["block_size"])
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_bad_entropy if=uge")
    _load(program, 2, FLAC_STATE["output_base"])
    _store(program, 2, FLAC_STATE["reconstruct_output"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_warmup_done if=eq")
    program.op("loop_setup 1, r1")
    program.label("flac_warmup_loop")
    _load(program, 1, FLAC_STATE["coded_bits"])
    program.op("call_fwd flac_get_signed")
    _load(program, 2, FLAC_STATE["reconstruct_output"])
    program.op("st32 r2, r15, 0")
    program.op("movi r3, 4")
    program.op("add r2, r2, r3")
    _store(program, 2, FLAC_STATE["reconstruct_output"])
    program.op("loop_back 1, flac_warmup_loop")
    program.label("flac_warmup_done")
    program.op("call_fwd flac_prepare_predictor")
    _load(program, 1, FLAC_STATE["kind"])
    program.op("movi r2, 32")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_residual_header if=ult")
    program.op("movi r1, 4")
    program.op("call_fwd flac_get_bits")
    program.op("movi r2, 1")
    program.op("add r15, r15, r2")
    program.op("movi r2, 16")
    program.op("cmp r15, r2")
    program.op("jump_fwd flac_bad_entropy if=eq")
    _store(program, 15, FLAC_STATE["parameter_bits"])
    program.op("movi r1, 5")
    program.op("call_fwd flac_get_signed")
    program.op("st32 r0, r15, 0x4208")
    _load(program, 2, FLAC_STATE["order"])
    program.op("movi r3, 0x4280")
    _store(program, 3, FLAC_STATE["coefficient_pointer"])
    program.op("loop_setup 1, r2")
    program.label("flac_lpc_coefficient_loop")
    _load(program, 1, FLAC_STATE["parameter_bits"])
    program.op("call_fwd flac_get_signed")
    _load(program, 3, FLAC_STATE["coefficient_pointer"])
    program.op("st32 r3, r15, 0")
    program.op("movi r4, 4")
    program.op("add r3, r3, r4")
    _store(program, 3, FLAC_STATE["coefficient_pointer"])
    program.op("loop_back 1, flac_lpc_coefficient_loop")
    program.label("flac_residual_header")
    program.op("movi r1, 2")
    program.op("call_fwd flac_get_bits")
    program.op("movi r1, 2")
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_bad_entropy if=uge")
    program.op("movi r1, 1")
    program.op("add r15, r15, r1")
    program.op("movi r1, 3")
    program.op("add r15, r15, r1")
    _store(program, 15, FLAC_STATE["parameter_bits"])
    program.op("movi r1, 4")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["partition_order"])
    program.op("movi r1, 1")
    program.op("shl r2, r1, r15")
    _load(program, 3, FLAC_STATE["block_size"])
    program.op("movi r4, 1")
    program.op("sub r4, r2, r4")
    program.op("and r4, r3, r4")
    program.op("cmp r4, r0")
    program.op("jump_fwd flac_bad_partition if=ne")
    _store(program, 0, FLAC_STATE["partition_index"])
    _store(program, 0, FLAC_STATE["residual_buffered"])
    program.op("loop_setup 1, r2")
    program.label("flac_partition_loop")
    _load(program, 1, FLAC_STATE["parameter_bits"])
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["rice_parameter"])
    program.op("movi r1, 1")
    _load(program, 2, FLAC_STATE["parameter_bits"])
    program.op("shl r1, r1, r2")
    program.op("movi r2, 1")
    program.op("sub r1, r1, r2")
    program.op("cmp r15, r1")
    program.op("jump_fwd flac_rice_normal if=ne")
    program.op("movi r1, 5")
    program.op("call_fwd flac_get_bits")
    _store(program, 15, FLAC_STATE["raw_width"])
    program.op("jump_fwd flac_partition_count")
    program.label("flac_rice_normal")
    program.op("movi r1, 0xff")
    _store(program, 1, FLAC_STATE["raw_width"])
    program.label("flac_partition_count")
    _load(program, 1, FLAC_STATE["block_size"])
    _load(program, 2, FLAC_STATE["partition_order"])
    program.op("shr r1, r1, r2")
    _load(program, 2, FLAC_STATE["partition_index"])
    program.op("cmp r2, r0")
    program.op("jump_fwd flac_partition_count_ready if=ne")
    _load(program, 2, FLAC_STATE["order"])
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_bad_partition if=ult")
    program.op("sub r1, r1, r2")
    program.label("flac_partition_count_ready")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_partition_empty if=eq")
    program.op("loop_setup 2, r1")
    program.label("flac_residual_loop")
    _load(program, 1, FLAC_STATE["raw_width"])
    program.op("movi r2, 0xff")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_residual_rice if=eq")
    program.op("cmp r1, r0")
    program.op("movi r15, 0 if=eq")
    program.op("jump_fwd flac_residual_ready if=eq")
    program.op("call_fwd flac_get_signed")
    program.op("jump_fwd flac_residual_ready")
    program.label("flac_residual_rice")
    _load(program, 1, FLAC_STATE["rice_parameter"])
    program.op("call_fwd flac_get_rice")
    program.label("flac_residual_ready")
    _load(program, 1, FLAC_STATE["residual_buffered"])
    program.op("movi r2, 2")
    program.op("shl r2, r1, r2")
    program.op("movi r3, 0x4100")
    program.op("add r2, r2, r3")
    program.op("st32 r2, r15, 0")
    program.op("movi r2, 1")
    program.op("add r1, r1, r2")
    _store(program, 1, FLAC_STATE["residual_buffered"])
    program.op("movi r2, 64")
    program.op("cmp r1, r2")
    program.op("call_fwd flac_reconstruct if=eq")
    program.op("loop_back 2, flac_residual_loop")
    program.label("flac_partition_empty")
    _load(program, 1, FLAC_STATE["partition_index"])
    program.op("movi r2, 1")
    program.op("add r1, r1, r2")
    _store(program, 1, FLAC_STATE["partition_index"])
    program.op("loop_back 1, flac_partition_loop")
    _load(program, 1, FLAC_STATE["residual_buffered"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_subframe_done if=eq")
    program.op("call_fwd flac_reconstruct_tail")
    program.label("flac_subframe_done")
    program.op("ret")

    program.label("flac_prepare_predictor")
    program.op("movi r1, 0x4280")
    program.op("movi r2, 0x4300")
    program.op("movi r3, 32")
    program.op("loop_setup 1, r3")
    program.label("flac_predictor_zero_loop")
    program.op("st32 r1, r0, 0")
    program.op("st32 r2, r0, 0")
    program.op("movi r4, 4")
    program.op("add r1, r1, r4")
    program.op("add r2, r2, r4")
    program.op("loop_back 1, flac_predictor_zero_loop")
    _load(program, 1, FLAC_STATE["order"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_prepare_params if=eq")
    _load(program, 2, FLAC_STATE["output_base"])
    program.op("movi r3, 2")
    program.op("shl r3, r1, r3")
    program.op("add r2, r2, r3")
    program.op("movi r3, 0x4300")
    program.op("loop_setup 1, r1")
    program.label("flac_history_loop")
    program.op("movi r4, 4")
    program.op("sub r2, r2, r4")
    program.op("ld32 r5, r2, 0")
    program.op("st32 r3, r5, 0")
    program.op("add r3, r3, r4")
    program.op("loop_back 1, flac_history_loop")
    program.label("flac_prepare_params")
    program.op("movi r1, 0x4280")
    program.op("st32 r0, r1, 0x4200")
    program.op("movi r1, 0x4300")
    program.op("st32 r0, r1, 0x4204")
    program.op("st32 r0, r0, 0x4208")
    _load(program, 1, FLAC_STATE["kind"])
    program.op("movi r2, 32")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_prepare_predictor_return if=uge")
    _load(program, 1, FLAC_STATE["order"])
    for order in range(1, 5):
        program.op(f"movi r2, {order}")
        program.op("cmp r1, r2")
        program.op(f"jump_fwd flac_fixed_coeff_{order} if=eq")
    program.op("jump_fwd flac_prepare_predictor_return")
    fixed_coefficients = {
        1: (1,),
        2: (2, -1),
        3: (3, -3, 1),
        4: (4, -6, 4, -1),
    }
    for order, coefficients in fixed_coefficients.items():
        program.label(f"flac_fixed_coeff_{order}")
        for index, value in enumerate(coefficients):
            program.op(f"movi r2, {value & 0xFFFFFFFF}")
            program.op(f"st32 r0, r2, 0x{0x4280 + index * 4:04x}")
        program.op("jump_fwd flac_prepare_predictor_return")
    program.op("nop")
    program.label("flac_prepare_predictor_return")
    program.op("ret")

    program.label("flac_reconstruct")
    _store(program, 0, FLAC_STATE["residual_buffered"])
    program.op("movi r1, 0x4100")
    program.op("movi r2, 0x4200")
    _load(program, 5, FLAC_STATE["reconstruct_output"])
    program.op("lpc r5, r1, r2, 32, 64")
    _load(program, 1, FLAC_STATE["reconstruct_output"])
    program.op("movi r2, 256")
    program.op("add r1, r1, r2")
    _store(program, 1, FLAC_STATE["reconstruct_output"])
    program.op("ret")

    program.label("flac_reconstruct_tail")
    program.op("movi r1, 0x4100")
    program.op("movi r2, 0x4200")
    _load(program, 7, FLAC_STATE["residual_buffered"])
    for count in (32, 16, 8, 4, 2, 1):
        skip = f"flac_reconstruct_tail_skip_{count}"
        program.op(f"movi r4, {count}")
        program.op("cmp r7, r4")
        program.op(f"jump_fwd {skip} if=ult")
        _load(program, 5, FLAC_STATE["reconstruct_output"])
        program.op(f"lpc r5, r1, r2, 32, {count}")
        program.op(f"movi r3, {count * 4}")
        program.op("add r1, r1, r3")
        _load(program, 5, FLAC_STATE["reconstruct_output"])
        program.op("add r5, r5, r3")
        _store(program, 5, FLAC_STATE["reconstruct_output"])
        program.op("sub r7, r7, r4")
        program.label(skip)
    _store(program, 0, FLAC_STATE["residual_buffered"])
    program.op("ret")

    program.label("flac_process_frame")
    _load(program, 1, FLAC_STATE["frame_input"])
    _load(program, 2, WAV_STATE["uncommitted_input"])
    program.op("add r2, r2, r1")
    _store(program, 2, WAV_STATE["uncommitted_input"])
    _store(program, 0, FLAC_STATE["frame_input"])
    _load(program, 1, WAV_STATE["profile"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_process_resampled if=ne")
    program.label("flac_process_passthrough")
    _load(program, 1, FLAC_STATE["block_size"])
    _store(program, 1, FLAC_STATE["frame_remaining"])
    _store(program, 0, FLAC_STATE["process_index"])
    program.op("movi r3, 128")
    program.op("loop_setup 1, r3")
    program.label("flac_output_tile_loop")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_output_tile_idle if=eq")
    program.op("movi r3, 32")
    program.op("min r3, r1, r3")
    _store(program, 3, WAV_STATE["tile_frames"])
    _store(program, 3, WAV_STATE["source_tile_frames"])
    program.op("movi r2, 0x4000")
    program.op("loop_setup 2, r3")
    program.label("flac_process_sample_loop")
    _load(program, 1, FLAC_STATE["process_index"])
    program.op("movi r3, 2")
    program.op("shl r3, r1, r3")
    program.op("ld32 r4, r3, 0")
    _load(program, 7, FLAC_STATE["wasted0"])
    program.op("shl r4, r4, r7")
    _load(program, 5, WAV_STATE["channels"])
    program.op("movi r6, 1")
    program.op("cmp r5, r6")
    program.op("jump_fwd flac_process_channels_ready if=eq")
    _load(program, 5, FLAC_STATE["block_size"])
    program.op("add r5, r5, r1")
    program.op("movi r6, 2")
    program.op("shl r5, r5, r6")
    program.op("ld32 r5, r5, 0")
    _load(program, 7, FLAC_STATE["wasted1"])
    program.op("shl r5, r5, r7")
    _load(program, 6, FLAC_STATE["assignment"])
    program.op("movi r7, 8")
    program.op("cmp r6, r7")
    program.op("sub r5, r4, r5 if=eq")
    program.op("movi r7, 9")
    program.op("cmp r6, r7")
    program.op("add r4, r4, r5 if=eq")
    program.op("movi r7, 10")
    program.op("cmp r6, r7")
    program.op("jump_fwd flac_process_channels_ready if=ne")
    program.op("movi r7, 1")
    program.op("and r6, r5, r7")
    program.op("shl r4, r4, r7")
    program.op("or r4, r4, r6")
    program.op("add r6, r4, r5")
    program.op("sub r5, r4, r5")
    program.op("sar r4, r6, r7")
    program.op("sar r5, r5, r7")
    program.label("flac_process_channels_ready")
    _load(program, 6, WAV_STATE["downmix"])
    program.op("cmp r6, r0")
    program.op("jump_fwd flac_process_no_downmix if=eq")
    program.op("movi r6, 1")
    program.op("sar r7, r4, r6")
    program.op("sar r8, r5, r6")
    program.op("add r7, r7, r8")
    program.op("and r8, r4, r6")
    program.op("and r6, r5, r6")
    program.op("add r8, r8, r6")
    program.op("movi r6, 1")
    program.op("cmp r8, r6")
    program.op("and r6, r7, r6")
    program.op("add r7, r7, r6 if=eq")
    program.op("mov r4, r7")
    program.op("jump_fwd flac_process_normalize")
    program.label("flac_process_no_downmix")
    _load(program, 6, WAV_STATE["channels"])
    _load(program, 7, WAV_STATE["output_channels"])
    program.op("movi r8, 1")
    program.op("cmp r6, r8")
    program.op("jump_fwd flac_process_normalize if=ne")
    program.op("movi r8, 2")
    program.op("cmp r7, r8")
    program.op("mov r5, r4 if=eq")
    program.label("flac_process_normalize")
    _load(program, 6, WAV_STATE["bits"])
    program.op("movi r7, 32")
    program.op("sub r6, r7, r6")
    program.op("shl r4, r4, r6")
    program.op("shl r5, r5, r6")
    program.op("st32 r2, r4, 0")
    _load(program, 6, WAV_STATE["output_channels"])
    program.op("movi r7, 1")
    program.op("cmp r6, r7")
    program.op("jump_fwd flac_process_advance if=eq")
    program.op("st32 r2, r5, 4")
    program.op("movi r7, 4")
    program.op("add r2, r2, r7")
    program.label("flac_process_advance")
    program.op("movi r7, 4")
    program.op("add r2, r2, r7")
    _load(program, 1, FLAC_STATE["process_index"])
    program.op("movi r7, 1")
    program.op("add r1, r1, r7")
    _store(program, 1, FLAC_STATE["process_index"])
    program.op("loop_back 2, flac_process_sample_loop")
    program.op("movi r2, 0x4000")
    _store(program, 2, WAV_STATE["pack_input"])
    program.op("movi r2, 0x4200")
    _store(program, 2, WAV_STATE["pack_output"])
    program.op("movi r4, 32")
    _store(program, 4, WAV_STATE["pack_capacity"])
    program.op("call_fwd flac_pack_tile")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    _load(program, 3, WAV_STATE["source_tile_frames"])
    program.op("sub r1, r1, r3")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    program.label("flac_output_tile_idle")
    program.op("loop_back 1, flac_output_tile_loop")
    program.op("ret")

    program.label("flac_process_resampled")
    # Rebuild a 25-frame interleaved window at 0x4000. The retained 16-frame
    # prefix contains seven history frames plus at most nine CRC-verified,
    # deliberately unissued tail frames from the preceding FLAC frame.
    program.op("movi r1, 0x4000")
    program.op("movi r2, 0")
    program.op("movi r3, 50")
    program.op("loop_setup 2, r3")
    program.label("flac_resample_window_clear")
    program.op("st32 r1, r2, 0")
    program.op("movi r4, 4")
    program.op("add r1, r1, r4")
    program.op("loop_back 2, flac_resample_window_clear")
    program.op("movi r1, 0x4680")
    program.op("movi r2, 0x4000")
    _load(program, 3, WAV_STATE["output_channels"])
    program.op("movi r4, 4")
    program.op("shl r3, r3, r4")
    program.op("loop_setup 2, r3")
    program.label("flac_resample_prefix_copy")
    program.op("ld32 r5, r1, 0")
    program.op("st32 r2, r5, 0")
    program.op("movi r4, 4")
    program.op("add r1, r1, r4")
    program.op("add r2, r2, r4")
    program.op("loop_back 2, flac_resample_prefix_copy")
    _store(program, 0, FLAC_STATE["process_index"])
    _load(program, 1, WAV_STATE["resample_pending"])
    program.op("movi r2, 7")
    program.op("add r2, r2, r1")
    program.op("movi r3, 2")
    program.op("shl r2, r2, r3")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("movi r5, 2")
    program.op("cmp r4, r5")
    program.op("movi r5, 1")
    program.op("shl r2, r2, r5 if=eq")
    program.op("movi r3, 0x4000")
    program.op("add r2, r2, r3")
    _store(program, 2, WAV_STATE["normalize_dst"])
    program.op("movi r2, 18")
    program.op("sub r2, r2, r1")
    _load(program, 3, FLAC_STATE["block_size"])
    program.op("min r2, r2, r3")
    program.op("loop_setup 2, r2")
    program.label("flac_resample_initial_fill")
    program.op("call_fwd flac_normalize_sample")
    program.op("loop_back 2, flac_resample_initial_fill")
    _load(program, 1, WAV_STATE["resample_pending"])
    _load(program, 2, FLAC_STATE["process_index"])
    program.op("add r1, r1, r2")
    _store(program, 1, WAV_STATE["resample_queued"])
    _load(program, 2, FLAC_STATE["block_size"])
    _load(program, 3, WAV_STATE["resample_pending"])
    program.op("add r2, r2, r3")
    program.op("loop_setup 1, r2")
    program.label("flac_resample_group_loop")
    _load(program, 1, WAV_STATE["resample_queued"])
    _load(program, 2, WAV_STATE["resample_final"])
    program.op("cmp r2, r0")
    program.op("jump_fwd flac_resample_final_target if=ne")
    program.op("movi r3, 10")
    program.op("cmp r1, r3")
    program.op("jump_fwd flac_resample_group_idle if=ult")
    program.op("movi r3, 9")
    program.op("sub r1, r1, r3")
    program.label("flac_resample_final_target")
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_resample_group_idle if=eq")
    program.op("movi r3, 8")
    program.op("min r1, r1, r3")
    _store(program, 0, WAV_STATE["tile_frames"])
    program.op("movi r1, 0x4100")
    _store(program, 1, WAV_STATE["resample_output"])
    _load(program, 1, WAV_STATE["resample_queued"])
    _load(program, 2, WAV_STATE["resample_final"])
    program.op("cmp r2, r0")
    program.op("movi r3, 9")
    program.op("sub r1, r1, r3 if=eq")
    program.op("movi r3, 8")
    program.op("min r1, r1, r3")
    program.op("loop_setup 2, r1")
    program.label("flac_resample_source_loop")
    _load(program, 3, WAV_STATE["output_channels"])
    program.op("movi r4, 2")
    program.op("shl r3, r3, r4")
    program.op("movi r1, 0x4000")
    program.op("movi r4, 7")
    program.op("loop_setup 3, r4")
    program.label("flac_resample_input_pointer")
    program.op("add r1, r1, r3")
    program.op("loop_back 3, flac_resample_input_pointer")
    program.op("movi r2, 0x4500")
    _load(program, 4, WAV_STATE["resample_output"])
    program.op("mov r8, r4")
    program.op("call_fwd flac_resample_one")
    _load(program, 5, WAV_STATE["tile_frames"])
    program.op("add r5, r5, r8")
    _store(program, 5, WAV_STATE["tile_frames"])
    _load(program, 6, WAV_STATE["output_channels"])
    program.op("movi r7, 2")
    program.op("shl r6, r6, r7")
    program.op("shl r7, r8, r7")
    program.op("movi r5, 2")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("cmp r4, r5")
    program.op("movi r5, 1")
    program.op("shl r7, r7, r5 if=eq")
    _load(program, 4, WAV_STATE["resample_output"])
    program.op("add r4, r4, r7")
    _store(program, 4, WAV_STATE["resample_output"])
    program.op("movi r1, 0x4000")
    program.op("add r2, r1, r6")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("movi r5, 24")
    program.op("movi r7, 2")
    program.op("cmp r4, r7")
    program.op("movi r7, 1")
    program.op("shl r5, r5, r7 if=eq")
    program.op("loop_setup 3, r5")
    program.label("flac_resample_window_shift")
    program.op("ld32 r7, r2, 0")
    program.op("st32 r1, r7, 0")
    program.op("movi r7, 4")
    program.op("add r1, r1, r7")
    program.op("add r2, r2, r7")
    program.op("loop_back 3, flac_resample_window_shift")
    _load(program, 1, FLAC_STATE["process_index"])
    _load(program, 2, FLAC_STATE["block_size"])
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_resample_append_zero if=uge")
    _load(program, 3, WAV_STATE["output_channels"])
    program.op("mov r4, r3")
    program.op("movi r5, 5")
    program.op("shl r3, r3, r5")
    program.op("movi r5, 6")
    program.op("shl r4, r4, r5")
    program.op("add r3, r3, r4")
    program.op("movi r4, 0x4000")
    program.op("add r3, r3, r4")
    _store(program, 3, WAV_STATE["normalize_dst"])
    program.op("call_fwd flac_normalize_sample")
    program.op("jump_fwd flac_resample_source_next")
    program.label("flac_resample_append_zero")
    _load(program, 1, WAV_STATE["output_channels"])
    program.op("mov r4, r1")
    program.op("mov r2, r1")
    program.op("movi r3, 5")
    program.op("shl r1, r1, r3")
    program.op("movi r3, 6")
    program.op("shl r2, r2, r3")
    program.op("add r2, r2, r1")
    program.op("movi r3, 0x4000")
    program.op("add r2, r2, r3")
    program.op("st32 r2, r0, 0")
    program.op("movi r3, 2")
    program.op("cmp r4, r3")
    program.op("st32 r2, r0, 4 if=eq")
    _load(program, 1, WAV_STATE["resample_queued"])
    program.op("movi r2, 1")
    program.op("sub r1, r1, r2")
    _store(program, 1, WAV_STATE["resample_queued"])
    program.label("flac_resample_source_next")
    program.op("loop_back 2, flac_resample_source_loop")
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_resample_group_idle if=eq")
    program.op("movi r2, 0x4100")
    _store(program, 2, WAV_STATE["pack_input"])
    _store(program, 2, WAV_STATE["pack_output"])
    program.op("call_fwd flac_pack_tile")
    program.label("flac_resample_group_idle")
    program.op("loop_back 1, flac_resample_group_loop")
    program.op("movi r1, 0x4000")
    program.op("movi r2, 0x4680")
    _load(program, 3, WAV_STATE["output_channels"])
    program.op("movi r4, 4")
    program.op("shl r3, r3, r4")
    program.op("loop_setup 2, r3")
    program.label("flac_resample_prefix_save")
    program.op("ld32 r5, r1, 0")
    program.op("st32 r2, r5, 0")
    program.op("movi r4, 4")
    program.op("add r1, r1, r4")
    program.op("add r2, r2, r4")
    program.op("loop_back 2, flac_resample_prefix_save")
    _load(program, 1, WAV_STATE["resample_queued"])
    _store(program, 1, WAV_STATE["resample_pending"])
    program.op("ret")

    program.label("flac_normalize_sample")
    _load(program, 1, FLAC_STATE["process_index"])
    program.op("movi r3, 2")
    program.op("shl r3, r1, r3")
    program.op("ld32 r4, r3, 0")
    _load(program, 7, FLAC_STATE["wasted0"])
    program.op("shl r4, r4, r7")
    _load(program, 5, WAV_STATE["channels"])
    program.op("movi r6, 1")
    program.op("cmp r5, r6")
    program.op("jump_fwd flac_normalize_channels_ready if=eq")
    _load(program, 5, FLAC_STATE["block_size"])
    program.op("add r5, r5, r1")
    program.op("movi r6, 2")
    program.op("shl r5, r5, r6")
    program.op("ld32 r5, r5, 0")
    _load(program, 7, FLAC_STATE["wasted1"])
    program.op("shl r5, r5, r7")
    _load(program, 6, FLAC_STATE["assignment"])
    program.op("movi r7, 8")
    program.op("cmp r6, r7")
    program.op("sub r5, r4, r5 if=eq")
    program.op("movi r7, 9")
    program.op("cmp r6, r7")
    program.op("add r4, r4, r5 if=eq")
    program.op("movi r7, 10")
    program.op("cmp r6, r7")
    program.op("jump_fwd flac_normalize_channels_ready if=ne")
    program.op("movi r7, 1")
    program.op("and r6, r5, r7")
    program.op("shl r4, r4, r7")
    program.op("or r4, r4, r6")
    program.op("add r6, r4, r5")
    program.op("sub r5, r4, r5")
    program.op("sar r4, r6, r7")
    program.op("sar r5, r5, r7")
    program.label("flac_normalize_channels_ready")
    _load(program, 6, WAV_STATE["downmix"])
    program.op("cmp r6, r0")
    program.op("jump_fwd flac_normalize_no_downmix if=eq")
    program.op("movi r6, 1")
    program.op("sar r7, r4, r6")
    program.op("sar r8, r5, r6")
    program.op("add r7, r7, r8")
    program.op("and r8, r4, r6")
    program.op("and r6, r5, r6")
    program.op("add r8, r8, r6")
    program.op("movi r6, 1")
    program.op("cmp r8, r6")
    program.op("and r6, r7, r6")
    program.op("add r7, r7, r6 if=eq")
    program.op("mov r4, r7")
    program.op("jump_fwd flac_normalize_store")
    program.label("flac_normalize_no_downmix")
    _load(program, 6, WAV_STATE["channels"])
    _load(program, 7, WAV_STATE["output_channels"])
    program.op("movi r8, 1")
    program.op("cmp r6, r8")
    program.op("jump_fwd flac_normalize_store if=ne")
    program.op("movi r8, 2")
    program.op("cmp r7, r8")
    program.op("mov r5, r4 if=eq")
    program.label("flac_normalize_store")
    _load(program, 6, WAV_STATE["bits"])
    program.op("movi r7, 32")
    program.op("sub r6, r7, r6")
    program.op("shl r4, r4, r6")
    program.op("shl r5, r5, r6")
    _load(program, 2, WAV_STATE["normalize_dst"])
    program.op("st32 r2, r4, 0")
    _load(program, 6, WAV_STATE["output_channels"])
    program.op("movi r7, 1")
    program.op("cmp r6, r7")
    program.op("jump_fwd flac_normalize_advance if=eq")
    program.op("st32 r2, r5, 4")
    program.op("movi r7, 4")
    program.op("add r2, r2, r7")
    program.label("flac_normalize_advance")
    program.op("movi r7, 4")
    program.op("add r2, r2, r7")
    _store(program, 2, WAV_STATE["normalize_dst"])
    _load(program, 1, FLAC_STATE["process_index"])
    program.op("movi r7, 1")
    program.op("add r1, r1, r7")
    _store(program, 1, FLAC_STATE["process_index"])
    program.op("ret")

    program.label("flac_resample_one")
    _load(program, 3, WAV_STATE["profile"])
    for profile in range(1, 16):
        program.op(f"movi r6, {profile}")
        program.op("cmp r3, r6")
        program.op(f"jump_fwd flac_resample_one_{profile} if=eq")
    program.op("jump_fwd flac_output_unsupported")
    for profile in range(1, 16):
        program.label(f"flac_resample_one_{profile}")
        program.op(f"resample r8, r1, r2, {profile}, 1")
        program.op("ret")

    program.label("flac_align_zero")
    program.op("cmp r9, r0")
    program.op("jump_fwd flac_align_return if=eq")
    program.op("mov r1, r9")
    program.op("call_fwd flac_get_bits")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_bad_pad if=ne")
    program.label("flac_align_return")
    program.op("ret")

    program.label("flac_get_signed")
    program.op("mov r5, r1")
    _store(program, 5, FLAC_STATE["signed_width"])
    program.op("call_fwd flac_get_bits")
    _load(program, 5, FLAC_STATE["signed_width"])
    program.op("movi r1, 32")
    program.op("sub r1, r1, r5")
    program.op("shl r15, r15, r1")
    program.op("sar r15, r15, r1")
    program.op("ret")

    program.label("flac_get_rice")
    program.op("mov r2, r1")
    program.op("movi r1, 0")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    _store(program, 1, FLAC_STATE["unary_done"])
    for index in range(32):
        program.op("call_fwd flac_get_bit")
        program.op("cmp r15, r0")
        program.op(f"jump_fwd flac_rice_found_{index} if=ne")
        _load(program, 1, FLAC_STATE["frame_remaining"])
        program.op("movi r3, 1")
        program.op("add r1, r1, r3")
        _store(program, 1, FLAC_STATE["frame_remaining"])
    # Preserve the 32-bit fast path. A long token first drains the partial byte,
    # then scans at most 8192 whole bytes so the 65535 quotient bound does not
    # consume one retired microinstruction per encoded bit.
    program.op("movi r1, 8")
    program.op("loop_setup 3, r1")
    program.label("flac_rice_partial_loop")
    _load(program, 1, FLAC_STATE["unary_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_rice_partial_idle if=ne")
    program.op("cmp r9, r0")
    program.op("jump_fwd flac_rice_partial_idle if=eq")
    program.op("call_fwd flac_get_bit")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_rice_partial_zero if=eq")
    program.op("movi r1, 1")
    _store(program, 1, FLAC_STATE["unary_done"])
    program.op("jump_fwd flac_rice_partial_idle")
    program.label("flac_rice_partial_zero")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("movi r3, 1")
    program.op("add r1, r1, r3")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    program.label("flac_rice_partial_idle")
    program.op("loop_back 3, flac_rice_partial_loop")
    _load(program, 1, FLAC_STATE["unary_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_rice_found if=ne")
    program.op("movi r1, 256")
    program.op("loop_setup 3, r1")
    program.label("flac_rice_chunk_loop")
    _load(program, 1, FLAC_STATE["unary_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_rice_chunk_idle if=ne")
    for index in range(32):
        _load(program, 1, FLAC_STATE["unary_done"])
        program.op("cmp r1, r0")
        program.op(f"jump_fwd flac_rice_chunk_byte_{index}_done if=ne")
        program.op("call_fwd flac_rice_scan_byte")
        program.label(f"flac_rice_chunk_byte_{index}_done")
    program.label("flac_rice_chunk_idle")
    program.op("loop_back 3, flac_rice_chunk_loop")
    _load(program, 1, FLAC_STATE["unary_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_rice_found if=ne")
    program.op("jump_fwd flac_profile_limit5")
    for index in range(32):
        program.label(f"flac_rice_found_{index}")
    program.label("flac_rice_found")
    program.op("mov r1, r2")
    program.op("cmp r1, r0")
    program.op("movi r15, 0 if=eq")
    program.op("jump_fwd flac_rice_value_ready if=eq")
    program.op("call_fwd flac_get_bits")
    program.label("flac_rice_value_ready")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    _load(program, 2, FLAC_STATE["rice_parameter"])
    program.op("shl r1, r1, r2")
    program.op("or r1, r1, r15")
    program.op("movi r2, 1")
    program.op("and r3, r1, r2")
    program.op("shr r15, r1, r2")
    program.op("sub r3, r0, r3")
    program.op("xor r15, r15, r3")
    _store(program, 0, FLAC_STATE["unary_done"])
    program.op("ret")

    program.label("flac_get_bits")
    program.op("movi r2, 0")
    program.op("loop_setup 3, r1")
    program.label("flac_get_bits_loop")
    program.op("call_fwd flac_get_bit")
    program.op("movi r3, 1")
    program.op("shl r2, r2, r3")
    program.op("or r2, r2, r15")
    program.op("loop_back 3, flac_get_bits_loop")
    program.op("mov r15, r2")
    program.op("ret")

    program.label("flac_get_unary")
    program.op("movi r2, 0")
    program.op("movi r1, 24")
    program.op("loop_setup 3, r1")
    program.label("flac_unary_loop")
    _load(program, 1, FLAC_STATE["unary_done"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_unary_idle if=ne")
    program.op("call_fwd flac_get_bit")
    program.op("cmp r15, r0")
    program.op("jump_fwd flac_unary_zero if=eq")
    program.op("movi r1, 1")
    _store(program, 1, FLAC_STATE["unary_done"])
    program.op("jump_fwd flac_unary_idle")
    program.label("flac_unary_zero")
    program.op("movi r1, 1")
    program.op("add r2, r2, r1")
    program.label("flac_unary_idle")
    program.op("loop_back 3, flac_unary_loop")
    program.op("mov r15, r2")
    _store(program, 0, FLAC_STATE["unary_done"])
    program.op("ret")

    program.label("flac_get_bit")
    program.op("cmp r9, r0")
    program.op("jump_fwd flac_have_bit_byte if=ne")
    program.op("cmp r13, r0")
    program.op("jump_fwd flac_have_fifo_word if=ne")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_have_batch if=ne")
    _load(program, 3, WAV_STATE["input_length"])
    program.op("cmp r11, r3")
    program.op("jump_fwd flac_truncated if=uge")
    program.op(f"movi r3, 0x{READER_STAGING:04x}")
    program.op("movi r4, 128")
    program.op("input_refill r14, r3, r4")
    program.op("dma_wait r5")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_truncated if=eq")
    program.label("flac_have_batch")
    program.op("fifo_pop r12")
    program.op("movi r3, 7")
    program.op("and r13, r13, r3")
    program.op("sub r14, r14, r13")
    program.label("flac_have_fifo_word")
    program.op("movi r3, 0xff")
    program.op("and r10, r12, r3")
    program.op("movi r3, 8")
    program.op("shr r12, r12, r3")
    program.op("movi r3, 1")
    program.op("sub r13, r13, r3")
    program.op("add r11, r11, r3")
    _store(program, 11, CTX_PARSED)
    program.op("and r3, r8, r3")
    program.op("cmp r3, r0")
    program.op("crc8 r7, r7, r10 if=ne")
    program.op("movi r3, 2")
    program.op("and r3, r8, r3")
    program.op("cmp r3, r0")
    program.op("crc16 r6, r6, r10 if=ne")
    program.op("movi r9, 8")
    program.label("flac_have_bit_byte")
    program.op("movi r3, 7")
    program.op("shr r15, r10, r3")
    program.op("movi r3, 1")
    program.op("and r15, r15, r3")
    program.op("shl r10, r10, r3")
    program.op("sub r9, r9, r3")
    program.op("ret")

    program.label("flac_rice_scan_byte")
    program.op("call_fwd flac_get_aligned_byte")
    program.op("mov r4, r15")
    program.op("cmp r4, r0")
    program.op("jump_fwd flac_rice_scan_zero if=eq")
    for index in range(8):
        program.op(f"movi r3, {0x80 >> index}")
        program.op("and r5, r4, r3")
        program.op("cmp r5, r0")
        program.op(f"jump_fwd flac_rice_scan_found_{index} if=ne")
    program.op("jump_fwd flac_profile_limit5")
    for index in range(8):
        program.label(f"flac_rice_scan_found_{index}")
        _load(program, 1, FLAC_STATE["frame_remaining"])
        if index != 0:
            program.op(f"movi r3, {index}")
            program.op("add r1, r1, r3")
        program.op("movi r3, 65535")
        program.op("cmp r3, r1")
        program.op("jump_fwd flac_profile_limit5 if=ult")
        _store(program, 1, FLAC_STATE["frame_remaining"])
        program.op("mov r10, r4")
        program.op(f"movi r3, {index + 1}")
        program.op("shl r10, r10, r3")
        program.op(f"movi r9, {7 - index}")
        program.op("movi r1, 1")
        _store(program, 1, FLAC_STATE["unary_done"])
        program.op("jump_fwd flac_rice_scan_return")
    program.label("flac_rice_scan_zero")
    _load(program, 1, FLAC_STATE["frame_remaining"])
    program.op("movi r3, 8")
    program.op("add r1, r1, r3")
    program.op("movi r3, 65535")
    program.op("cmp r3, r1")
    program.op("jump_fwd flac_profile_limit5 if=ult")
    _store(program, 1, FLAC_STATE["frame_remaining"])
    program.label("flac_rice_scan_return")
    program.op("ret")

    program.label("flac_get_aligned_byte")
    program.op("cmp r9, r0")
    program.op("jump_fwd flac_bad_entropy if=ne")
    program.op("cmp r13, r0")
    program.op("jump_fwd flac_aligned_have_word if=ne")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_aligned_have_batch if=ne")
    _load(program, 3, WAV_STATE["input_length"])
    program.op("cmp r11, r3")
    program.op("jump_fwd flac_truncated if=uge")
    program.op(f"movi r3, 0x{READER_STAGING:04x}")
    program.op("movi r4, 128")
    program.op("input_refill r14, r3, r4")
    program.op("dma_wait r5")
    program.op("cmp r14, r0")
    program.op("jump_fwd flac_truncated if=eq")
    program.label("flac_aligned_have_batch")
    program.op("fifo_pop r12")
    program.op("movi r3, 7")
    program.op("and r13, r13, r3")
    program.op("sub r14, r14, r13")
    program.label("flac_aligned_have_word")
    program.op("movi r3, 0xff")
    program.op("and r15, r12, r3")
    program.op("movi r3, 8")
    program.op("shr r12, r12, r3")
    program.op("movi r3, 1")
    program.op("sub r13, r13, r3")
    program.op("add r11, r11, r3")
    _store(program, 11, CTX_PARSED)
    program.op("and r3, r8, r3")
    program.op("cmp r3, r0")
    program.op("crc8 r7, r7, r15 if=ne")
    program.op("movi r3, 2")
    program.op("and r3, r8, r3")
    program.op("cmp r3, r0")
    program.op("crc16 r6, r6, r15 if=ne")
    program.op("ret")

    # FLAC duplicates the two output helpers because APUMC calls are forward-only.
    program.label("flac_pack_tile")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_pack_preflush_done if=eq")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("movi r3, 256")
    program.op("cmp r2, r3")
    program.op("call_fwd flac_emit_pending if=eq")
    _load(program, 1, WAV_STATE["output_mode"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_pack_preflush_done if=ne")
    _load(program, 1, CTX_OUTPUT_CAPACITY)
    _load(program, 2, CTX_OUTPUT_CURSOR)
    program.op("sub r1, r1, r2")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("cmp r1, r2")
    program.op("call_fwd flac_emit_pending if=eq")
    program.label("flac_pack_preflush_done")
    _load(program, 1, WAV_STATE["output_mode"])
    program.op("cmp r1, r0")
    program.op("call_fwd p5_limit_pack if=eq")
    _load(program, 1, WAV_STATE["pending_input"])
    _load(program, 2, WAV_STATE["uncommitted_input"])
    program.op("add r1, r1, r2")
    _store(program, 1, WAV_STATE["pending_input"])
    _store(program, 0, WAV_STATE["uncommitted_input"])
    _load(program, 1, WAV_STATE["tile_frames"])
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 1")
    program.op("cmp r2, r3")
    program.op("shl r1, r1, r3 if=ne")
    _load(program, 3, WAV_STATE["pack_input"])
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r6, 16")
    program.op("movi r7, 8")
    program.op("cmp r4, r6")
    program.op("mov r7, r6 if=eq")
    program.op("loop_setup 2, r1")
    program.label("flac_pack_shift_loop")
    program.op("ld32 r5, r3, 0")
    program.op("sar r5, r5, r7")
    program.op("st32 r3, r5, 0")
    program.op("movi r6, 4")
    program.op("add r3, r3, r6")
    program.op("loop_back 2, flac_pack_shift_loop")
    program.op("movi r3, 0x40000000")
    program.op("st32 r0, r3, 0x4540")
    program.op("st32 r0, r2, 0x4544")
    program.op("movi r1, 0x4570")
    _store(program, 1, WAV_STATE["pending_output"])
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("loop_setup 3, r1")
    program.label("flac_pack_frame_loop")
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_pack_frame_idle if=eq")
    _load(program, 6, WAV_STATE["pending_bytes"])
    program.op("movi r7, 256")
    program.op("cmp r6, r7")
    program.op("call_fwd flac_emit_pending if=eq")
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("movi r6, 8")
    program.op("min r6, r1, r6")
    program.op("movi r7, 256")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("sub r7, r7, r2")
    program.op("movi r3, 1")
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r5, 16")
    program.op("cmp r4, r5")
    program.op("movi r5, 1")
    program.op("add r3, r3, r5 if=ne")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("cmp r4, r5")
    program.op("add r3, r3, r5 if=ne")
    program.op("shr r7, r7, r3")
    program.op("min r6, r6, r7")
    _store(program, 6, WAV_STATE["pack_capacity"])
    _load(program, 1, WAV_STATE["pack_input"])
    program.op("movi r2, 0x4540")
    program.op("movi r5, 0x4570")
    _load(program, 7, WAV_STATE["pending_bytes"])
    program.op("add r5, r5, r7")
    _load(program, 4, WAV_STATE["output_bits"])
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("call_fwd p5_pack_small")
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 2")
    program.op("shl r2, r2, r3")
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("movi r7, 0")
    program.op("loop_setup 2, r6")
    program.label("flac_pack_input_advance_loop")
    program.op("add r7, r7, r2")
    program.op("loop_back 2, flac_pack_input_advance_loop")
    program.op("add r1, r1, r7")
    _store(program, 1, WAV_STATE["pack_input"])
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r4, r7")
    program.op("movi r4, 2")
    program.op("movi r7, 4")
    program.op("mov r4, r7 if=ne")
    _load(program, 2, WAV_STATE["output_channels"])
    program.op("movi r3, 1")
    program.op("cmp r2, r3")
    program.op("shl r4, r4, r3 if=ne")
    _load(program, 6, WAV_STATE["pack_capacity"])
    program.op("movi r7, 0")
    program.op("loop_setup 2, r6")
    program.label("flac_pack_byte_count_loop")
    program.op("add r7, r7, r4")
    program.op("loop_back 2, flac_pack_byte_count_loop")
    _load(program, 2, WAV_STATE["pending_bytes"])
    program.op("add r2, r2, r7")
    _store(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 2, WAV_STATE["pending_frames"])
    program.op("add r2, r2, r6")
    _store(program, 2, WAV_STATE["pending_frames"])
    _load(program, 1, WAV_STATE["tile_frames"])
    program.op("sub r1, r1, r6")
    _store(program, 1, WAV_STATE["tile_frames"])
    program.op("movi r3, 1")
    _store(program, 3, WAV_STATE["pending"])
    program.label("flac_pack_frame_idle")
    program.op("loop_back 3, flac_pack_frame_loop")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("call_fwd flac_emit_pending if=ne")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("jump_fwd p5_capacity_fail if=ne")
    program.op("ret")

    program.label("p5_limit_pack")
    _load(program, 1, CTX_OUTPUT_CAPACITY)
    _load(program, 2, CTX_OUTPUT_CURSOR)
    program.op("cmp r1, r2")
    program.op("jump_fwd p5_capacity_fail if=ult")
    program.op("sub r1, r1, r2")
    _load(program, 6, WAV_STATE["pending_bytes"])
    program.op("cmp r1, r6")
    program.op("jump_fwd p5_capacity_fail if=ult")
    program.op("sub r1, r1, r6")
    program.op("movi r3, 1")
    _load(program, 4, WAV_STATE["output_bits"])
    program.op("movi r5, 16")
    program.op("cmp r4, r5")
    program.op("movi r5, 1")
    program.op("add r3, r3, r5 if=ne")
    _load(program, 4, WAV_STATE["output_channels"])
    program.op("cmp r4, r5")
    program.op("add r3, r3, r5 if=ne")
    program.op("shr r1, r1, r3")
    _load(program, 2, WAV_STATE["tile_frames"])
    program.op("min r1, r1, r2")
    program.op("cmp r1, r0")
    program.op("jump_fwd p5_capacity_fail if=eq")
    program.op("cmp r1, r2")
    program.op("movi r3, 1")
    program.op(f"st32 r0, r3, 0x{WAV_STATE['capacity_stop']:04x} if=ult")
    _store(program, 1, WAV_STATE["tile_frames"])
    program.label("p5_limit_return")
    program.op("ret")

    program.label("p5_pack_small")
    for count in range(1, 9):
        program.op(f"movi r3, {count}")
        program.op("cmp r6, r3")
        program.op(f"jump_fwd p5_pack_small_{count} if=eq")
    program.op("jump_fwd flac_output_unsupported")
    for count in range(1, 9):
        program.label(f"p5_pack_small_{count}")
        program.op("movi r7, 16")
        program.op("cmp r4, r7")
        program.op(f"pcm_pack r5, r1, r2, 0, {count} if=eq")
        program.op(f"pcm_pack r5, r1, r2, 1, {count} if=ne")
        program.op("ret")

    program.label("flac_emit_pending")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_emit_return if=eq")
    _load(program, 1, WAV_STATE["pending_output"])
    _load(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 3, WAV_STATE["output_mode"])
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_emit_stream if=ne")
    program.op("output_commit r8, r1, r2")
    program.op("jump_fwd flac_emit_commit")
    program.label("flac_emit_stream")
    _load(program, 3, WAV_STATE["inflight"])
    program.op("cmp r3, r0")
    program.op("call_fwd flac_retire_inflight if=ne")
    _load(program, 1, WAV_STATE["pending_output"])
    _load(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 5, CTX_FINAL)
    _store(program, 5, WAV_STATE["inflight_final"])
    program.op("output_stream r10, r1, r2")
    _load(program, 3, WAV_STATE["pending_frames"])
    _store(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["pending_input"])
    _store(program, 4, WAV_STATE["inflight_input"])
    program.op("movi r3, 1")
    _store(program, 3, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["pending"])
    _store(program, 0, WAV_STATE["pending_bytes"])
    _store(program, 0, WAV_STATE["pending_frames"])
    _store(program, 0, WAV_STATE["pending_input"])
    _load(program, 1, WAV_STATE["inflight_final"])
    program.op("cmp r1, r0")
    program.op("call_fwd flac_retire_inflight if=ne")
    program.op("jump_fwd flac_emit_capacity")
    program.label("flac_emit_commit")
    program.op("dma_wait r8")
    _load(program, 3, WAV_STATE["pending_frames"])
    _load(program, 4, WAV_STATE["pending_input"])
    _load(program, 5, WAV_STATE["output_channels"])
    _load(program, 6, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r6, r7")
    program.op("movi r6, 2")
    program.op("jump_fwd flac_emit_sample_bytes if=eq")
    program.op("movi r6, 4")
    program.label("flac_emit_sample_bytes")
    program.op("movi r7, 1")
    program.op("cmp r5, r7")
    program.op("jump_fwd flac_emit_frame_bytes if=eq")
    program.op("shl r6, r6, r7")
    program.label("flac_emit_frame_bytes")
    program.op("loop_setup 2, r3")
    program.label("flac_emit_commit_loop")
    program.op("frame_commit r4, r6")
    program.op("movi r4, 0")
    program.op("loop_back 2, flac_emit_commit_loop")
    _store(program, 0, WAV_STATE["pending"])
    _store(program, 0, WAV_STATE["pending_bytes"])
    _store(program, 0, WAV_STATE["pending_frames"])
    _store(program, 0, WAV_STATE["pending_input"])
    program.label("flac_emit_capacity")
    _load(program, 1, WAV_STATE["capacity_stop"])
    program.op("cmp r1, r0")
    program.op("jump_fwd p5_capacity_fail if=ne")
    program.label("flac_emit_return")
    program.op("ret")

    program.label("flac_retire_inflight")
    _load(program, 1, WAV_STATE["inflight"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_retire_return if=eq")
    program.op("dma_wait r8")
    _load(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["inflight_input"])
    _load(program, 5, WAV_STATE["output_channels"])
    _load(program, 6, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r6, r7")
    program.op("movi r6, 2")
    program.op("jump_fwd flac_retire_sample_bytes if=eq")
    program.op("movi r6, 4")
    program.label("flac_retire_sample_bytes")
    program.op("movi r7, 1")
    program.op("cmp r5, r7")
    program.op("jump_fwd flac_retire_frame_bytes if=eq")
    program.op("shl r6, r6, r7")
    program.label("flac_retire_frame_bytes")
    program.op("loop_setup 2, r3")
    program.label("flac_retire_commit_loop")
    program.op("frame_commit r4, r6")
    program.op("movi r4, 0")
    program.op("loop_back 2, flac_retire_commit_loop")
    _store(program, 0, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["inflight_frames"])
    _store(program, 0, WAV_STATE["inflight_input"])
    _store(program, 0, WAV_STATE["inflight_final"])
    program.label("flac_retire_return")
    program.op("ret")

    # Exact format diagnostics. Warnings are retained on later failures.
    for label, code, reason in (
        ("flac_bad_signature", 4, 0x0010),
        ("flac_bad_streaminfo", 4, 0x0015),
        ("flac_bad_metadata", 4, 0x0016),
        ("flac_bad_frame", 4, 0x0017),
        ("flac_geometry_change", 4, 0x0018),
        ("flac_bad_pad", 4, 0x0019),
        ("flac_bad_number", 4, 0x001A),
        ("flac_sample_count", 4, 0x0032),
        ("flac_bad_entropy", 7, 0x0040),
        ("flac_bad_wasted", 7, 0x0041),
        ("flac_bad_partition", 7, 0x0042),
        ("flac_header_crc", 6, 0x0030),
        ("flac_frame_crc", 6, 0x0031),
        ("flac_unsupported_encoding", 3, 0x0003),
        ("flac_unsupported_geometry", 3, 0x0002),
        ("flac_profile_limit4", 3, 0x0004),
    ):
        program.label(label)
        program.op(f"movi r4, {code}")
        program.op(f"movi r5, 0x{reason:04x}")
        if label in ("flac_geometry_change", "flac_sample_count"):
            _store(program, 4, WAV_STATE["error_code"])
            _store(program, 5, WAV_STATE["error_reason"])
            program.op("call_fwd flac_error_drain")
            _load(program, 4, WAV_STATE["error_code"])
            _load(program, 5, WAV_STATE["error_reason"])
        program.op("jump_fwd flac_fail4")
    program.label("flac_profile_limit5")
    program.op("movi r4, 3")
    program.op("movi r5, 0x0004")
    program.op("jump_fwd flac_fail5")
    program.label("flac_output_unsupported")
    program.op("movi r4, 3")
    program.op("movi r5, 0x0050")
    program.op("jump_fwd flac_fail0")
    program.label("flac_truncated")
    program.op("movi r4, 5")
    program.op("movi r5, 0x0020")
    _load(program, 3, WAV_STATE["input_length"])
    _store(program, 3, CTX_DIAGNOSTIC)
    _load(program, 1, FLAC_STATE["parsing_stage"])
    program.op("movi r2, 5")
    program.op("cmp r1, r2")
    program.op("jump_fwd flac_fail5 if=eq")
    program.op("jump_fwd flac_fail4")

    # Geometry/count failures occur between frames, so all packed prefix data is
    # authoritative. Drain it without marking a failed stream command final.
    program.label("flac_error_drain")
    _store(program, 0, CTX_FINAL)
    program.op("call_fwd flac_error_retire")
    program.op("call_fwd flac_error_emit")
    program.op("call_fwd flac_error_retire")
    program.op("ret")

    program.label("flac_error_emit")
    _load(program, 1, WAV_STATE["pending"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_error_emit_return if=eq")
    _load(program, 1, WAV_STATE["pending_output"])
    _load(program, 2, WAV_STATE["pending_bytes"])
    _load(program, 3, WAV_STATE["output_mode"])
    program.op("cmp r3, r0")
    program.op("jump_fwd flac_error_emit_stream if=ne")
    program.op("output_commit r8, r1, r2")
    program.op("jump_fwd flac_error_emit_record")
    program.label("flac_error_emit_stream")
    program.op("output_stream r10, r1, r2")
    program.label("flac_error_emit_record")
    _load(program, 3, WAV_STATE["pending_frames"])
    _store(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["pending_input"])
    _store(program, 4, WAV_STATE["inflight_input"])
    program.op("movi r3, 1")
    _store(program, 3, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["pending"])
    _store(program, 0, WAV_STATE["pending_bytes"])
    _store(program, 0, WAV_STATE["pending_frames"])
    _store(program, 0, WAV_STATE["pending_input"])
    program.label("flac_error_emit_return")
    program.op("ret")

    program.label("flac_error_retire")
    _load(program, 1, WAV_STATE["inflight"])
    program.op("cmp r1, r0")
    program.op("jump_fwd flac_error_retire_return if=eq")
    program.op("dma_wait r8")
    _load(program, 3, WAV_STATE["inflight_frames"])
    _load(program, 4, WAV_STATE["inflight_input"])
    _load(program, 5, WAV_STATE["output_channels"])
    _load(program, 6, WAV_STATE["output_bits"])
    program.op("movi r7, 16")
    program.op("cmp r6, r7")
    program.op("movi r6, 2")
    program.op("jump_fwd flac_error_retire_sample_bytes if=eq")
    program.op("movi r6, 4")
    program.label("flac_error_retire_sample_bytes")
    program.op("movi r7, 1")
    program.op("cmp r5, r7")
    program.op("jump_fwd flac_error_retire_frame_bytes if=eq")
    program.op("shl r6, r6, r7")
    program.label("flac_error_retire_frame_bytes")
    program.op("loop_setup 2, r3")
    program.label("flac_error_retire_commit_loop")
    program.op("frame_commit r4, r6")
    program.op("movi r4, 0")
    program.op("loop_back 2, flac_error_retire_commit_loop")
    _store(program, 0, WAV_STATE["inflight"])
    _store(program, 0, WAV_STATE["inflight_frames"])
    _store(program, 0, WAV_STATE["inflight_input"])
    _store(program, 0, WAV_STATE["inflight_final"])
    program.label("flac_error_retire_return")
    program.op("ret")

    for stage in (0, 4, 5, 6):
        program.label(f"flac_fail{stage}")
        _load(program, 1, WAV_STATE["warnings"])
        program.op("movi r2, 16")
        program.op("shl r1, r1, r2")
        program.op("or r5, r5, r1")
        program.op("movi r1, 2")
        program.op("movi r2, 24")
        program.op("shl r1, r1, r2")
        program.op("or r5, r5, r1")
        program.op(f"job_result r4, r5, {stage}")
        program.op("end")
    program.label("p5_capacity_fail")
    program.op("movi r4, 8")
    program.op("movi r5, 0x0051")
    _load(program, 1, WAV_STATE["warnings"])
    program.op("movi r2, 16")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    _load(program, 1, WAV_STATE["format"])
    program.op("movi r2, 24")
    program.op("shl r1, r1, r2")
    program.op("or r5, r5, r1")
    program.op("job_result r4, r5, 8")
    program.op("end")


def microcode_source() -> str:
    program = Program()
    program.line("# Generated by scripts/generate_apu_p5_microcode.py; edit the emitter.")
    program.line(".build_id 0x0005000000000002")
    program.line(".entry 0 wav wav wav_last 65535 16777215 0x1800 0x4800 0x001ff83b 0 0x1800")
    program.line(".entry 1 mp3 mp3 mp3 1 1 0 0 0 0 0")
    program.line(".entry 2 flac flac flac_last 65535 16777215 0x1800 0x4800 0x001ff83b 0 0x1800")
    program.line()
    program.label("mp3")
    program.op("trap 0x01000001")
    program.line()
    wav_program = Program()
    _emit_wav(wav_program)
    helper_index = wav_program.lines.index("wav_validate_alignment:")
    program.lines.extend(wav_program.lines[:helper_index])
    program.line()
    _emit_flac(program)
    program.line()
    program.lines.extend(wav_program.lines[helper_index:])
    program.line()
    _emit_flac_helpers(program)
    program.label("wav_last")
    program.op("end")
    program.label("flac_last")
    program.op("end")
    return program.text()


def main() -> int:
    parser = argparse.ArgumentParser(description=__doc__)
    parser.add_argument("--output", type=Path, required=True)
    parser.add_argument("--check", action="store_true")
    args = parser.parse_args()
    expected = microcode_source()
    if args.check:
        try:
            actual = args.output.read_text(encoding="utf-8")
        except OSError:
            return 1
        return 0 if actual == expected else 1
    args.output.write_text(expected, encoding="utf-8")
    return 0


if __name__ == "__main__":
    raise SystemExit(main())