Squashed from work; the fine-grained history is under tag archive/work-2026-08-26:
- vj: the tweener learns — RIFE runs on the Mac and feeds the warp
- asset-ai: the chat tells the truth while it works
- llm: a cold turn on the solo slot takes the session-native speculative path — think-mode turns no longer re-ingest the whole conversation through the draft head (66 → 122 tok/s on the four-lane box)
- llm: the step cost model is chosen per device — the RTX PRO 6000's measured verify curve (13.7 + 3.17·B ms) beside the 5090's; the bench warms every tail shape and times two windows
- vj: local store, lyrics and model plumbing, and the frame-interpolator's device parity check
A Turing box (RTX 2080 Ti, sm_75) lost ALL of CUDA because two kernel
files refused to compile for it, and one failed kernel build means the
stub store — surfaced in the VJ as "stems: model error: no compiled-graph
device" on the DJ tab.
diffusion_ops.cu used three sm_80-only pieces unguarded: bf16 wmma
fragments (the type itself is incomplete before Ampere), cp.async, and
the m16n8k16 mma shapes. The cp.async helpers now fall back to
synchronous copies below sm_80 — the f16 wmma flash/sdpa kernels lose
their prefetch overlap on Turing, not their contents — while the bf16
and FA2 kernels are compiled out and their launchers refuse pre-sm_80
devices with cudaErrorNotSupported instead of returning a buffer the
kernel never wrote.
fattn/common.cuh made mkllm_unused_vars constexpr: the no-cp.async
branch of ggml_cuda_fattn_mma_get_nstages calls it, and a non-constexpr
callee poisoned the constexpr config chain on exactly the pre-Ampere
device pass — the arch nobody had compiled for.
Stems verified on the 2080 Ti: stems-ops-check all green (SNR 137-147 dB
against the CPU reference), two tracks separated end to end, output
confirmed clean by ear.
Co-authored-by: vjroger <r.deleeuw@qogni.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
The model code was spread across eight crates that had grown into each other:
ggml and cuda and mlx each owned part of a tensor runtime, llama and tts and
voice2 each owned part of a model, and libs/diffusion owned everything else.
They are now one tree with an explicit shape:
libs/ai/cuda — kernels and launch surface
libs/ai/metal — Metal shaders and the shim
libs/ai/llm — the language-model runtime (sessions, lanes, contexts,
the CUDA and Metal executors, the compiled Metal path)
libs/ai/models/ — common, flux, h3, music, paint, speech, stems, vision
libs/diffusion is not deleted but demoted: what remains is the VALIDATOR
crate — several dozen `*_validate.rs` oracles that check a native
implementation against a reference, which is where they belong now that the
implementations live next door.
The functional work inside the move is mostly in the LLM runtime: N lanes that
draft while one verify batch serves all of them, per-slot prefill over a shared
folded attention arena, speculation that survives batching, and a scheduler
that reports rather than publishes. And in the CUDA build: a machine without
usable CUDA must still LINK (and say so), the default kernel arch is the
building machine's GPU, `NO_CUDA` forces the stub even where the toolkit
exists, and kernels compile in parallel with progress.
libs/video_flow is new here: classical optical flow estimation and the `mkfl`
motion-field payload — a flow field measured from a clip without a model,
which is what drives free-rate bounce-looping playback and the uprez/tween
enhance pipe.
