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>
413709b565