The FP8 backbone mode added each linear's bias by downloading, tiling and
re-uploading it per call, which cost more than the FP8 GEMM saved. A
gpu_add_cols_broadcast op (CUDA kernel; host loop on the Metal tensor
backend) adds a cols-wide bias to every row on the device, and each linear
keeps its bias resident after the first upload.
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
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.
2026-08-23 01:34:35 +02:00
Renamed from libs/ggml/src/backend/cuda/diffusion_ops.cu (Browse further)
