iris: cut test debug info, say which adapter drew, and log on the desktop

Three findings from one morning, all of them things that were invisible
rather than wrong. docs/RUST.md's two new sections have the full account.

**`cargo test --workspace` was taking half an hour, and it was debug
info.** rustc's default `debug = true`, times eight test binaries each
statically linking the whole wgpu + naga + winit + parley graph, means
every one of them gets a private copy of that graph's DWARF written into
it: the linkers for one run had written ~54 GB between them and were
still going at thirty minutes -- the worst single one 16.9 GB for one
test binary -- leaving an 88 GB target/. It was not CPU: the machine was
87% idle, and rust-lld's threads were in D state in btrfs
`handle_reserve_ticket`, blocked on space reservation at 83% full. So
`debug = "line-tables-only"` on both `profile.dev` and `profile.test` --
both, because `cargo test` builds dependencies under one and the test
targets under the other. Cold, with all 19 suites run: 69 s and a 3.7 GB
target. Backtraces keep file and line; `RUSTFLAGS="-C debuginfo=2"` per
run buys back variable inspection when a debugger actually needs it.

**The desktop had no logger at all**, so every `log::` call on that side
went to `log`'s no-op default -- including the GLES fallback warning
added hours earlier. `DefaultApp::run` installs a stderr logger
(`src/default/logging.rs`, no new dependency: a level and a line is a
page of code against env_logger plus its filter dialect), and the
renderer now says which adapter won at `info`. That line is the point:
with a silent fallback, a layer-2 screenshot rendered by llvmpipe and one
rendered by the host's GPU are the same PNG, and which one it was is
exactly what the screenshot is being taken to judge.

**`tests/mask_sdf.rs` is a render pass now, not a compute pass.** It
asked for `adapter.limits()` because `iris_core::device_limits()`
deliberately zeroes the six `max_compute_*` fields -- a decision on
record since 2026-09-05, which this quietly worked around instead of
following. It now asks for what iris asks for and calls the function from
the fragment stage, where the renderer calls it. The compute pass was
*not* why it crashed, and the record should not say it was: the rewrite
crashes identically. What the crash is: dropping a wgpu device on this
VM's Venus adapter segfaults, after the test has produced its answer
(worst CPU/shader disagreement 5.8e-6). Narrowed -- plain Vulkan creating
and destroying five VkDevices on the same adapter is clean, and the same
binary with Vulkan hidden falls back to GL and exits clean. Worked around
at `Gpu::leak`, with the reason and the delete-me condition written
there.

`rigs/virtgpu-probe` is the new rig behind the Venus half: which capsets
the host offers (0x16 -- VIRGL, VIRGL2, VENUS; no capset 6, so no DRM
native context without host-side work), whether the device has compute
(it does: 1024 invocations/workgroup -- the "no compute" finding on
record is about the Android emulator's SwiftShader, a different machine),
and whether plain Vulkan teardown is clean.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

docs/RUST.md

@@ -7440,6 +7440,183 @@ re-derived:
 7. Decisions belong here with a date and what was rejected, the way
    `PLAN.md` does it. Do not put design into commit messages alone.
 
+### `cargo test --workspace` was taking half an hour, and it was debug info (2026-09-08)
+
+Iris asked why. It is worth reading before concluding that this machine
+needs more disk, because that was the first guess here too and it is the
+smaller half.
+
+**What it was.** rustc's default for `dev` is `debug = true`, and this
+workspace has **eight test binaries**, each statically linking the whole
+wgpu + naga + winit + parley graph. Every one of them therefore gets its
+own full copy of that graph's DWARF written into it at link time.
+Measured: the linkers for one `cargo test --workspace` had written
+**~54 GB between them** and were still going at 30 minutes, the worst
+single one **16.9 GB for one test binary**, leaving an **88 GB**
+`target/`.
+
+**What it was not: CPU.** The machine was **87% idle** with 8-12%
+iowait throughout, and every `rustc` was in `__futex_wait` with no busy
+thread. An earlier guess in this same session -- that concurrent agents'
+builds were loading the machine -- was wrong, and checking `ps` for a
+busy process would have killed it in a minute. `rust-lld`'s worker
+threads were in **`D` state in btrfs `handle_reserve_ticket`**: blocked
+on space reservation, because the filesystem was at 83% full. So the
+write volume is the disease and the full disk is what turned slow into
+stalled -- they move about 20 MB/s between them in that state.
+
+**How to tell it apart from a busy build**, since the two look identical
+from the outside (`cargo` printing `Compiling ...` and not finishing):
+
+    ps -eo pid,stat,etime,pcpu,comm | grep -E 'rustc|rust-lld'
+    cat /proc/<lld-pid>/wchan          # handle_reserve_ticket = btrfs, not you
+    for l in $(pgrep -x rust-lld); do awk '/^write_bytes/{print $2/1048576}' /proc/$l/io; done
+
+A linker at 0.5% CPU having written gigabytes is not compiling.
+
+**The fix, in `iris/Cargo.toml`.** `debug = "line-tables-only"` on both
+`profile.dev` and `profile.test` -- both, because `cargo test` builds
+dependencies under `dev` and the test targets themselves under `test`, so
+setting only `dev` leaves the eight big binaries at the default. It keeps
+the file and line of every frame, which is what a panicking test prints
+and what gdb needs to name the frames of a segfault (both were used
+today); it gives up inspecting variables in a debugger, which is worth
+asking for per-run when it is wanted:
+
+    RUSTFLAGS="-C debuginfo=2" cargo test -p iris --test mask_sdf
+
+**Measured after**, cold (`cargo clean` first), in `iris/`: the whole
+workspace compiles *and runs all 19 test suites* in **69 s**, with a
+**3.7 GB** `target/`. Before, an *incremental* run — everything already
+compiled, only the eight binaries left to link — had not finished after
+thirty minutes. `cargo clean` alongside the profile change freed
+**104 GB** (83% -> 67% full).
+
+**One trap while measuring this, worth not repeating**: the session's
+working directory is the repo root, not `iris/`, so a `du -sh target`
+appended to a build command reports the root workspace's 29 MB and makes
+a correct `iris` build look like it built the wrong thing. Read the test
+*names* in the log to tell two workspaces apart, not a path-relative
+size.
+
+### Venus went away for an hour, and nothing said so (2026-09-08)
+
+**It came back on its own and nothing in the host config changed.** Iris
+asked whether something had, since she still passes Venus as true. What
+happened, and what was ruled out, so the next occurrence is not
+re-investigated from scratch:
+
+The symptom, around 02:10 on 2026-09-08 with the VM at load 68 and
+several agents building: `vulkaninfo` reported `Failed to detect any
+valid GPUs in the current config` and `vkEnumeratePhysicalDevices failed
+with ERROR_INITIALIZATION_FAILED`; Mesa printed **`No virgl contexts
+available on host`**; and `wgpu` reported `NotFound { active_backends:
+VULKAN, no_adapter_backends: VULKAN, supported_backends: VULKAN | GL }`.
+So **both** paths through the virtio-gpu died at once -- Venus for
+Vulkan and virgl for GL -- and GL then fell through to llvmpipe, which
+is what actually rendered that hour's layer-2 screenshots.
+
+That string is Mesa's virgl DRM winsys, next to `DRM_IOCTL_VIRTGPU_
+CONTEXT_INIT failed with %s` in `libgallium`: the *host* refused a new
+context. After a reboot the same host config gives
+`Virtio-GPU Venus (AMD Radeon RX 7900 XT (RADV NAVI31))`, Mesa 26.2.2,
+driverID `MESA_VENUS`.
+
+What was ruled out, by measurement rather than by reasoning:
+
+- **A guest-side context cap.** 48 concurrent short-lived Vulkan clients
+  all succeed, and 120 concurrent *long-lived* ones (each holding a
+  `VkDevice` open at once, a throwaway holder) all succeed. So the
+  ceiling, if there is one, is not near the handful of GPU-using
+  processes that were running.
+- **A Mesa upgrade.** `mesa 1:26.1.7 -> 1:26.2.2` landed 2026-09-05, two
+  days before Venus was last seen working here.
+- **Anything in iris.** It was `vulkaninfo`'s answer too, from a
+  process that has never linked against this repo.
+
+**What this host's virtio-gpu actually offers**, asked of the kernel
+rather than assumed (`rigs/virtgpu-probe`): bitmask `0x16` --
+**VIRGL, VIRGL2 and VENUS**. Capset 6, the **DRM "native context"**, is
+not offered. That is the answer to "is there something to do with qemu
+instead of Venus", asked by Iris 2026-09-08: native context is the thing
+worth wanting -- RADV running *in the guest* against a passed-through DRM
+context instead of Venus proxying every Vulkan call, and it is where
+Mesa's effort has gone. Whether it would avoid the teardown crash below
+is **untested** -- it is a different driver stack, so it is a reasonable
+thing to try rather than a known fix -- but it needs the **host** side to
+offer it (virglrenderer built with
+its amdgpu DRM renderer, and a qemu that exposes `context_types=drm`;
+crosvm has it further along). The guest would also need `vulkan-radeon`
+installed, which it does not have today -- only `vulkan-virtio`. The
+other two options are VFIO passthrough (complete, but the host loses the
+GPU) and dropping `venus=true` (leaves virgl/GL only, i.e. no Vulkan at
+all, which is the wrong direction since Vulkan is the phone's path).
+
+So it is host-side and transient, and this VM cannot see the host to say
+more: no `dmesg` (the guest's kernel buffer is not readable to this
+user), and no view of the host's `amdgpu`. **The honest state is "we do
+not know which host-side resource ran out"** -- worth capturing the host
+side of it if it recurs, since that is the half that would answer it.
+
+**A second, unrelated Venus fault, found the moment it came back
+(2026-09-08).** `iris/tests/mask_sdf.rs` had passed the day before and
+now `SIGSEGV`d -- and it had passed *because Venus was down*, so it
+silently ran on GL. What it actually is, narrowed by measurement:
+
+- **The test's work completes and its answer is right** (worst
+  CPU/shader disagreement 5.8e-6). The crash is at process teardown,
+  dropping wgpu's device: a call through an unmapped address on a
+  wgpu-created thread, per gdb.
+- **It is wgpu's teardown, not Venus's device lifecycle.** A plain
+  Vulkan program creating and destroying five `VkDevice`s and its
+  instance on the same adapter exits cleanly (`rigs/virtgpu-probe`).
+- **It is Venus-specific.** The same test binary, with Vulkan hidden
+  (`VK_DRIVER_FILES=/nonexistent`) so wgpu falls back to GL, exits
+  cleanly.
+
+Worked around in the test rather than fixed, at `Gpu::leak` with the
+reason written there: one device for the whole test, handed to the
+process instead of dropped. **Compute was investigated and is not
+involved** -- an early version of that test used a compute pass, which
+was wrong for its own reason (the paragraph after this one), but the
+render-pass rewrite crashes identically, and Venus here reports full
+compute anyway (`maxComputeWorkGroupInvocations` 1024,
+`maxComputeSharedMemorySize` 65536, Vulkan 1.4 -- `rigs/virtgpu-probe`
+again). The 2026-09-05 "no compute" finding is about the **Android
+emulator's SwiftShader GL path** reporting ES 3.0, which is a different
+machine; it says nothing about this VM. This was got wrong out loud
+first, so it is written down: the compute pass was blamed for the crash
+before the rewrite showed the crash was not about compute at all.
+
+The test was rewritten to a render pass regardless, and that part is not
+a workaround: it now asks for `iris_core::device_limits()` -- what iris
+itself requests -- and calls the function from the fragment stage, which
+is where the renderer calls it. A test that needs a capability the thing
+under test has never needed is testing the wrong device.
+
+**What was fixed, because the failure was silent.** Two things, both the
+rule that a degraded state must be distinguishable from a healthy one:
+
+1. `default::render::UiRenderer::new` had the defect the Android backend
+   was fixed for in `85869d0` -- `Backends::PRIMARY` and an `.expect` --
+   so layer 2 aborted with `Could not get adapter!` instead of falling
+   back. It now probes and rebuilds on `Backends::GL`, as Android does.
+2. **The desktop had no logger at all**, so every `log::` call on that
+   side -- including that new fallback warning -- went to `log`'s
+   no-op default. `DefaultApp::run` installs a stderr logger now
+   (`src/default/logging.rs`, no new dependency), and the renderer says
+   which adapter won at `info`:
+
+       INFO iris::default::render: iris renderer: Virtio-GPU Venus
+       (AMD Radeon RX 7900 XT (RADV NAVI31)) (Vulkan, venus Mesa
+       26.2.2-arch1.1) on Backends(VULKAN | METAL | DX12 | BROWSER_WEBGPU)
+
+   That line is the point: with the fallback in place and no logger, a
+   `run-headless.sh` screenshot rendered by llvmpipe and one rendered by
+   the host's GPU are the same PNG, and the difference is exactly what a
+   screenshot is being taken to judge. **Check it before trusting a
+   layer-2 screenshot or any frame number from that window.**
+
 ### Vulkan in the emulator (measured 2026-09-04)
 
 **Settled 2026-09-04: the guest gets Vulkan from SwiftShader, and the

iris/Cargo.toml

@@ -113,6 +113,35 @@ exclude = ["android-app"]
 version = "0.1.0"
 edition = "2024"
 
+# Debug info is the reason a `cargo test --workspace` here was taking half
+# an hour, and it is worth the paragraph. Measured 2026-09-08: with rustc's
+# default `debug = true`, linking this workspace's test binaries wrote
+# **~54 GB** (one single test binary's linker wrote 16.9 GB) and left an
+# **88 GB** `target/`. Eight test binaries each statically link the whole
+# wgpu + naga + winit + parley graph, and at the default every one of them
+# gets a full copy of that graph's DWARF written into it. On a btrfs at 83%
+# full the linkers then sat in `handle_reserve_ticket` -- uninterruptible,
+# waiting on space reservation -- at about 20 MB/s between them, which is
+# what "the tests are slow" actually was. Not CPU: the machine was 87% idle
+# throughout.
+#
+# `line-tables-only` keeps what is actually read from a backtrace -- the
+# file and line of every frame, which is what a panicking test prints and
+# what gdb needs to name the frames of a segfault. What it gives up is
+# inspecting variables in a debugger; when that is wanted, ask for it on
+# the command line for that one run rather than paying for it on every
+# build:
+#
+#     RUSTFLAGS="-C debuginfo=2" cargo test -p iris --test whatever
+[profile.dev]
+debug = "line-tables-only"
+
+# The tests are what this is really for; `cargo test` uses `dev` for
+# dependencies and `test` for the test targets themselves, so setting only
+# `dev` leaves the eight big binaries at the default.
+[profile.test]
+debug = "line-tables-only"
+
 [workspace.dependencies]
 pollster = "0.4.0"
 winit = "0.30.12"

iris/src/default/app.rs

@@ -27,6 +27,10 @@ pub struct App<State: AppState> {
 
 impl<State: AppState> App<State> {
     pub fn run() {
+        // The desktop's `main` in everything but name -- see
+        // `super::logging`'s doc for why the logger goes here and what
+        // its absence hid.
+        super::logging::install(log::LevelFilter::Info);
         let event_loop = EventLoop::with_user_event().build().unwrap();
         let proxy = event_loop.create_proxy();
         event_loop

iris/src/default/logging.rs

@@ -0,0 +1,86 @@
+//! A stderr logger for the desktop entry point.
+//!
+//! Without one, `log::` calls on this side go nowhere: `log`'s default is
+//! a no-op logger, and nothing in `desktop-app` or the examples ever
+//! installed a real one. That is how iris came to have a renderer that
+//! silently fell back to GLES (and, on this VM, on to llvmpipe when the
+//! host lost its virtio-gpu contexts) with **no record anywhere of what
+//! drew the frame** -- a layer-2 screenshot off llvmpipe and one off the
+//! host GPU are the same PNG, and the difference is exactly what a
+//! screenshot is being taken to judge.
+//!
+//! Installed by [`DefaultApp::run`](super::app::DefaultApp::run) rather
+//! than by a library call somewhere, because that function already takes
+//! over the process -- it owns the event loop and does not return -- so
+//! it is the desktop's `main` in everything but name, and one install
+//! there covers `desktop-app` and every example at once. `try_init`
+//! rather than `init`: a binary that installed its own logger first keeps
+//! it, and a second `DefaultApp::run` in one process is not an error.
+//!
+//! Deliberately not `env_logger`. All this owes the reader is a level and
+//! a line, which is a page of code against a dependency plus its own
+//! filter dialect; the Android side is `android_logger` for the same
+//! reason -- one line per platform's own convention.
+
+use std::io::Write;
+
+use log::{Level, LevelFilter, Log, Metadata, Record};
+
+/// Reads one level name from `RUST_LOG` -- `off`, `error`, `warn`,
+/// `info`, `debug`, `trace`, case-insensitively. **Not env_logger's
+/// per-module filter syntax**: anything else is ignored and the default
+/// stands, rather than being silently read as "off", since a typo that
+/// turned logging off would be indistinguishable from a quiet program.
+fn level_from_env(default: LevelFilter) -> LevelFilter {
+    match std::env::var("RUST_LOG") {
+        Ok(text) => text.trim().parse().unwrap_or(default),
+        Err(_) => default,
+    }
+}
+
+struct StderrLogger {
+    level: LevelFilter,
+}
+
+impl Log for StderrLogger {
+    fn enabled(&self, metadata: &Metadata) -> bool {
+        metadata.level() <= self.level
+    }
+
+    fn log(&self, record: &Record) {
+        if !self.enabled(record.metadata()) {
+            return;
+        }
+        // One write, not a `writeln!` per part: two threads logging at
+        // once interleave otherwise, and the frame and input traces are
+        // both written from whichever thread produced them.
+        let line = format!(
+            "{level:<5} {target}: {args}\n",
+            level = match record.level() {
+                Level::Error => "ERROR",
+                Level::Warn => "WARN",
+                Level::Info => "INFO",
+                Level::Debug => "DEBUG",
+                Level::Trace => "TRACE",
+            },
+            target = record.target(),
+            args = record.args(),
+        );
+        let _ = std::io::stderr().write_all(line.as_bytes());
+    }
+
+    fn flush(&self) {
+        let _ = std::io::stderr().flush();
+    }
+}
+
+/// Installs the stderr logger unless this process already has one.
+/// Defaults to `info`, which is where the renderer says which adapter it
+/// got; `RUST_LOG=debug` adds iris's own per-frame lines.
+pub fn install(default: LevelFilter) {
+    let level = level_from_env(default);
+    let logger = Box::leak(Box::new(StderrLogger { level }));
+    if log::set_logger(logger).is_ok() {
+        log::set_max_level(level);
+    }
+}

iris/src/default/mod.rs

@@ -15,6 +15,7 @@ mod access;
 mod app;
 mod attr;
 mod input;
+mod logging;
 mod platform;
 mod render;
 

iris/src/default/render.rs

@@ -132,6 +132,29 @@ impl UiRenderer {
                 panic!("No usable GPU adapter for backends {backends:?}: {error}")
             });
 
+        // Say which adapter won, the way the Android backend's own report
+        // does. Without it a layer-2 screenshot or frame time from this
+        // window carries no record of what drew it, and the two cases that
+        // matter look identical in the PNG: the host GPU through Venus,
+        // and llvmpipe after this VM lost its virtio-gpu contexts. That
+        // happened on 2026-09-08, and the only reason anyone noticed is
+        // that the fallback above did not exist yet and the app aborted
+        // instead. A silent fallback needs this line to stay honest.
+        {
+            let info = adapter.get_info();
+            log::info!(
+                "iris renderer: {name} ({backend:?}, {driver}{driver_info}) on {backends:?}",
+                name = info.name,
+                backend = info.backend,
+                driver = info.driver,
+                driver_info = if info.driver_info.is_empty() {
+                    String::new()
+                } else {
+                    format!(" {}", info.driver_info)
+                },
+            );
+        }
+
         // No features beyond what wgpu asks for by default, and no
         // binding-array limits: the atlas is one texture_2d_array and a
         // standalone image is its own ordinary bind group, neither of which

iris/tests/mask_sdf.rs

@@ -10,13 +10,28 @@
 //! or two off, which is exactly what nobody notices.
 //!
 //! So this runs **the real shader text**, lifted out of
-//! `iris_core::SHAPE_SHADER` by name rather than copied here, in a compute
+//! `iris_core::SHAPE_SHADER` by name rather than copied here, over a grid
-//! pass over a grid of points, and compares what came back with the Rust
+//! of points, and compares what came back with the Rust function at the
-//! function at the same points. This is the only test in the workspace
+//! same points. This is the only test in the workspace that needs a GPU;
-//! that needs a GPU; everything else about masks is layer 1 (docs/RUST.md's
+//! everything else about masks is layer 1 (docs/RUST.md's "Three test
-//! "Three test layers"). It fails rather than skips when there is no
+//! layers"). It fails rather than skips when there is no adapter, because
-//! adapter, because a check that quietly did not run reads exactly like a
+//! a check that quietly did not run reads exactly like a check that
-//! check that passed.
+//! passed.
+//!
+//! **It is a render pass, and it asks for `iris_core::device_limits()`,
+//! because those are the two things iris itself does.** The first version
+//! of this test was a compute pass, which meant asking for compute limits
+//! that `device_limits()` deliberately zeroes -- docs/RUST.md, 2026-09-05:
+//! nothing in `iris`/`iris-core` creates a `ComputePipeline` or writes a
+//! `@compute` stage, so the limits stopped being requested rather than a
+//! fallback being built for a capability nothing uses. A test that needs
+//! a capability the thing under test has never needed is testing the
+//! wrong device, which is reason enough.
+//!
+//! It is **not** why that version crashed, and the record should not say
+//! it was: the render-pass rewrite crashes in exactly the same place, and
+//! this VM's Venus reports full compute anyway. See `Gpu::leak` for what
+//! the crash actually is.
 
 use iris_core::{SHAPE_SHADER, rounded_rect_coverage, util::Vec2};
 use pollster::FutureExt;
@@ -35,6 +50,16 @@ const BOT_RIGHT: Vec2 = Vec2::new(170.75, 90.0);
 /// `min(edge, radius)` stops mattering.
 const RADII: [f32; 5] = [0.0, 0.75, 8.0, 20.0, 34.0];
 
+/// The grid, as an attachment: one texel per probe point. `GRID_W` is a
+/// multiple of 64 so that a row of `R32Float` is 256-byte aligned, which
+/// is what `copy_texture_to_buffer` requires; at `STEP` this spans the
+/// rect above and about four pixels of margin on every side, so the
+/// feather is sampled rather than stepped over.
+const GRID_W: u32 = 384;
+const GRID_H: u32 = 192;
+const STEP: f32 = 0.5;
+const ORIGIN: Vec2 = Vec2::new(TOP_LEFT.x - 4.0, TOP_LEFT.y - 4.0);
+
 /// f32 arithmetic in two compilers, not one: `length`/`sqrt` and
 /// `smoothstep` are each allowed a unit or two in the last place, and the
 /// GPU may contract a multiply-add the CPU does not. A coverage is in
@@ -45,18 +70,36 @@ const TOLERANCE: f32 = 1e-5;
 
 #[test]
 fn mask_sdf_matches_the_shader() {
-    let points = grid();
+    let gpu = Gpu::open();
-    let gpu = run_shader(&points);
     let mut worst = 0.0f32;
     let mut worst_at = (Vec2::new(0.0, 0.0), 0.0f32, 0.0f32, 0.0f32);
-    for (&(pos, radius), &got) in points.iter().zip(&gpu) {
+    let (mut inside, mut feather, mut outside) = (0u32, 0u32, 0u32);
-        let want = rounded_rect_coverage(pos, TOP_LEFT, BOT_RIGHT, radius);
+
-        let diff = (got - want).abs();
+    for radius in RADII {
-        if diff > worst {
+        let coverage = run_shader(&gpu, radius);
-            worst = diff;
+        for y in 0..GRID_H {
-            worst_at = (pos, radius, want, got);
+            for x in 0..GRID_W {
+                let pos = probe_at(x, y);
+                let got = coverage[(y * GRID_W + x) as usize];
+                let want = rounded_rect_coverage(pos, TOP_LEFT, BOT_RIGHT, radius);
+                let diff = (got - want).abs();
+                if diff > worst {
+                    worst = diff;
+                    worst_at = (pos, radius, want, got);
+                }
+                if got > 0.999 {
+                    inside += 1;
+                } else if got > 0.001 {
+                    feather += 1;
+                } else {
+                    outside += 1;
+                }
+            }
         }
     }
+
+    gpu.leak();
+
     let (pos, radius, want, got) = worst_at;
     assert!(
         worst <= TOLERANCE,
@@ -67,9 +110,6 @@ fn mask_sdf_matches_the_shader() {
     );
 
     // The half that would pass on a function returning a constant.
-    let inside = gpu.iter().filter(|c| **c > 0.999).count();
-    let feather = gpu.iter().filter(|c| **c > 0.001 && **c < 0.999).count();
-    let outside = gpu.iter().filter(|c| **c < 0.001).count();
     assert!(
         inside > 0 && feather > 0 && outside > 0,
         "the grid never crossed an edge ({inside} in, {feather} on the feather, {outside} out), \
@@ -77,81 +117,95 @@ fn mask_sdf_matches_the_shader() {
     );
 }
 
-/// Points spanning the rect and a margin outside it, at half-pixel steps
+/// The probe position of texel `(x, y)` -- the one place the mapping
-/// so the feather is sampled rather than stepped over, paired with each
+/// lives, so the CPU side and the fragment stage cannot walk different
-/// radius.
+/// grids.
-fn grid() -> Vec<(Vec2, f32)> {
+fn probe_at(x: u32, y: u32) -> Vec2 {
-    let mut points = Vec::new();
+    Vec2::new(ORIGIN.x + x as f32 * STEP, ORIGIN.y + y as f32 * STEP)
-    for radius in RADII {
-        let mut y = TOP_LEFT.y - 4.0;
-        while y <= BOT_RIGHT.y + 4.0 {
-            let mut x = TOP_LEFT.x - 4.0;
-            while x <= BOT_RIGHT.x + 4.0 {
-                points.push((Vec2::new(x, y), radius));
-                x += 0.5;
-            }
-            y += 0.5;
-        }
-    }
-    points
 }
 
-/// `shader.wgsl`'s own `rounded_rect_coverage`, evaluated at every point.
+/// `shader.wgsl`'s own `rounded_rect_coverage`, evaluated at every texel
-fn run_shader(points: &[(Vec2, f32)]) -> Vec<f32> {
+/// of an `R32Float` attachment: one fragment per grid point, read back
-    let instance = wgpu::Instance::default();
+/// whole. A fragment stage because that is the stage the function is
-    let adapter = instance
+/// really called from, so what this compares is the code path that draws
-        .request_adapter(&wgpu::RequestAdapterOptions::default())
+/// rather than a second one built to be measurable.
-        .block_on()
+fn run_shader(gpu: &Gpu, radius: f32) -> Vec<f32> {
-        .expect(
+    let Gpu { device, queue, .. } = gpu;
-            "no wgpu adapter on this machine, so the CPU/shader SDF agreement went unchecked. \
-             This VM has a virtio-gpu render node (see iris/run-headless.sh); if that is gone, \
-             fix it rather than deleting this test.",
-        );
-    let (device, queue) = adapter
-        .request_device(&wgpu::DeviceDescriptor {
-            // The adapter's own, not `iris_core::device_limits()`: those
-            // are what the *app* asks for, chosen down to what a phone
-            // GPU has, and they carry no compute at all
-            // (`max_compute_invocations_per_workgroup` is 0). This probe
-            // draws nothing and shares no pipeline with the app -- it
-            // runs one function to see what it returns.
-            required_limits: adapter.limits(),
-            ..Default::default()
-        })
-        .block_on()
-        .expect("could not get a device from the adapter");
-
     let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
         label: Some("mask sdf probe"),
         source: wgpu::ShaderSource::Wgsl(probe_source().into()),
     });
-    let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+
+    // R32Float, not an 8-bit colour format: a coverage quantised to 1/255
+    // could not be compared against the CPU's at anything like TOLERANCE,
+    // and the comparison would then be measuring the texture rather than
+    // the two functions.
+    let texture = device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("mask sdf coverage"),
+        size: wgpu::Extent3d {
+            width: GRID_W,
+            height: GRID_H,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::R32Float,
+        usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC,
+        view_formats: &[],
+    });
+    let view = texture.create_view(&Default::default());
+
+    let probe = Probe {
+        top_left: [TOP_LEFT.x, TOP_LEFT.y],
+        bot_right: [BOT_RIGHT.x, BOT_RIGHT.y],
+        origin: [ORIGIN.x, ORIGIN.y],
+        step: [STEP, STEP],
+        radius,
+        _pad: [0.0; 3],
+    };
+    let uniform = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
         label: Some("mask sdf probe"),
-        layout: None,
+        contents: bytemuck::bytes_of(&probe),
-        module: &module,
+        usage: wgpu::BufferUsages::UNIFORM,
-        entry_point: Some("probe"),
-        compilation_options: Default::default(),
-        cache: None,
     });
 
-    // (x, y, radius, unused) -- one vec4 per point, so the buffer needs no
+    let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
-    // stride arithmetic and no alignment rule of its own.
+        label: Some("mask sdf probe"),
-    let input: Vec<[f32; 4]> = points
+        layout: None,
-        .iter()
+        vertex: wgpu::VertexState {
-        .map(|(pos, radius)| [pos.x, pos.y, *radius, 0.0])
+            module: &module,
-        .collect();
+            entry_point: Some("probe_vs"),
-    let in_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            compilation_options: Default::default(),
-        label: Some("mask sdf points"),
+            buffers: &[],
-        contents: bytemuck::cast_slice(&input),
+        },
-        usage: wgpu::BufferUsages::STORAGE,
+        fragment: Some(wgpu::FragmentState {
+            module: &module,
+            entry_point: Some("probe_fs"),
+            compilation_options: Default::default(),
+            targets: &[Some(wgpu::TextureFormat::R32Float.into())],
+        }),
+        primitive: Default::default(),
+        depth_stencil: None,
+        multisample: Default::default(),
+        multiview_mask: None,
+        cache: None,
     });
-    let out_size = (points.len() * std::mem::size_of::<f32>()) as u64;
+    let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-    let out_buf = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("mask sdf probe"),
-        label: Some("mask sdf coverage"),
+        layout: &pipeline.get_bind_group_layout(0),
-        size: out_size,
+        entries: &[wgpu::BindGroupEntry {
-        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
+            binding: 0,
-        mapped_at_creation: false,
+            resource: uniform.as_entire_binding(),
+        }],
     });
+
+    // `copy_texture_to_buffer` wants each row 256-byte aligned; GRID_W is
+    // chosen so that it already is, rather than padding and unpicking the
+    // padding on the way out.
+    let row_bytes = GRID_W * 4;
+    assert_eq!(row_bytes % 256, 0, "GRID_W must keep rows 256-byte aligned");
+    let out_size = u64::from(row_bytes) * u64::from(GRID_H);
     let read_buf = device.create_buffer(&wgpu::BufferDescriptor {
         label: Some("mask sdf readback"),
         size: out_size,
@@ -159,29 +213,44 @@ fn run_shader(points: &[(Vec2, f32)]) -> Vec<f32> {
         mapped_at_creation: false,
     });
 
-    let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-        label: Some("mask sdf probe"),
-        layout: &pipeline.get_bind_group_layout(0),
-        entries: &[
-            wgpu::BindGroupEntry {
-                binding: 0,
-                resource: in_buf.as_entire_binding(),
-            },
-            wgpu::BindGroupEntry {
-                binding: 1,
-                resource: out_buf.as_entire_binding(),
-            },
-        ],
-    });
-
     let mut enc = device.create_command_encoder(&Default::default());
     {
-        let mut pass = enc.begin_compute_pass(&Default::default());
+        let mut pass = enc.begin_render_pass(&wgpu::RenderPassDescriptor {
+            label: Some("mask sdf probe"),
+            color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                view: &view,
+                depth_slice: None,
+                resolve_target: None,
+                ops: wgpu::Operations {
+                    load: wgpu::LoadOp::Clear(wgpu::Color::BLACK),
+                    store: wgpu::StoreOp::Store,
+                },
+            })],
+            depth_stencil_attachment: None,
+            multiview_mask: None,
+            timestamp_writes: None,
+            occlusion_query_set: None,
+        });
         pass.set_pipeline(&pipeline);
         pass.set_bind_group(0, &bind_group, &[]);
-        pass.dispatch_workgroups(points.len().div_ceil(64) as u32, 1, 1);
+        pass.draw(0..3, 0..1);
     }
-    enc.copy_buffer_to_buffer(&out_buf, 0, &read_buf, 0, out_size);
+    enc.copy_texture_to_buffer(
+        texture.as_image_copy(),
+        wgpu::TexelCopyBufferInfo {
+            buffer: &read_buf,
+            layout: wgpu::TexelCopyBufferLayout {
+                offset: 0,
+                bytes_per_row: Some(row_bytes),
+                rows_per_image: Some(GRID_H),
+            },
+        },
+        wgpu::Extent3d {
+            width: GRID_W,
+            height: GRID_H,
+            depth_or_array_layers: 1,
+        },
+    );
     queue.submit([enc.finish()]);
 
     let slice = read_buf.slice(..);
@@ -194,6 +263,89 @@ fn run_shader(points: &[(Vec2, f32)]) -> Vec<f32> {
     coverage
 }
 
+/// One device for the whole test, **never dropped** -- see [`Gpu::open`].
+struct Gpu {
+    instance: wgpu::Instance,
+    device: wgpu::Device,
+    queue: wgpu::Queue,
+}
+
+impl Gpu {
+    /// Opens the device this test draws with, and reports which adapter
+    /// answered, because that is not a detail here: this same test passes
+    /// on GL and **`SIGSEGV`s on Venus**, so a run that does not say which
+    /// one it got cannot be read.
+    fn open() -> Self {
+        let instance = wgpu::Instance::default();
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions::default())
+            .block_on()
+            .expect(
+                "no wgpu adapter on this machine, so the CPU/shader SDF agreement went \
+                 unchecked. This VM has a virtio-gpu render node (see iris/run-headless.sh); \
+                 if that is gone, fix it rather than deleting this test.",
+            );
+        let info = adapter.get_info();
+        eprintln!(
+            "mask_sdf: {} ({:?}, {})",
+            info.name, info.backend, info.driver
+        );
+        let (device, queue) = adapter
+            .request_device(&wgpu::DeviceDescriptor {
+                // What iris itself asks for -- see this file's header.
+                required_limits: iris_core::device_limits(),
+                ..Default::default()
+            })
+            .block_on()
+            .expect("could not get a device from the adapter");
+        Self {
+            instance,
+            device,
+            queue,
+        }
+    }
+
+    /// Hands the device to the process rather than dropping it.
+    ///
+    /// **The escape, and why there is no alternative here.** Dropping a
+    /// `wgpu` device against this VM's Venus adapter segfaults: a call
+    /// through an unmapped address on a wgpu-created thread, measured
+    /// 2026-09-08 with gdb. It is `wgpu`'s teardown and not the driver's
+    /// device lifecycle -- a plain Vulkan program creating and destroying
+    /// five `VkDevice`s and its instance on the same adapter exits
+    /// cleanly, and this same test binary exits cleanly when Vulkan is
+    /// hidden and it falls back to GL. Nothing this test can do about
+    /// wgpu's drop order makes that call valid, and the alternative is a
+    /// test that reports a crash after it has already produced its
+    /// answer, which is indistinguishable from the test failing.
+    ///
+    /// Safe because the process is about to exit: the leak is one device
+    /// and one instance, for the microseconds between here and `main`
+    /// returning. **Delete this the moment the teardown crash is fixed**
+    /// -- `cargo test -p iris --test mask_sdf` failing with `SIGSEGV`
+    /// after printing `test result: ok` is what it looks like when it is
+    /// still needed, and an ordinary pass is what it looks like when it
+    /// is not.
+    fn leak(self) {
+        std::mem::forget(self.queue);
+        std::mem::forget(self.device);
+        std::mem::forget(self.instance);
+    }
+}
+
+/// What the fragment stage needs to turn its own texel into a probe
+/// position: the rect being sampled, and where texel (0, 0) sits.
+#[repr(C)]
+#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+struct Probe {
+    top_left: [f32; 2],
+    bot_right: [f32; 2],
+    origin: [f32; 2],
+    step: [f32; 2],
+    radius: f32,
+    _pad: [f32; 3],
+}
+
 /// The probe module: the two functions **lifted from `shader.wgsl`
 /// itself**, plus an entry point that calls the outer one. Lifted rather
 /// than copied so there is nothing to keep in step -- an edit to the
@@ -202,21 +354,26 @@ fn run_shader(points: &[(Vec2, f32)]) -> Vec<f32> {
 fn probe_source() -> String {
     format!(
         "{}\n{}\n\
-         @group(0) @binding(0) var<storage, read> probe_in: array<vec4<f32>>;\n\
+         struct Probe {{\n\
-         @group(0) @binding(1) var<storage, read_write> probe_out: array<f32>;\n\
+         top_left: vec2<f32>,\n\
-         @compute @workgroup_size(64)\n\
+         bot_right: vec2<f32>,\n\
-         fn probe(@builtin(global_invocation_id) gid: vec3<u32>) {{\n\
+         origin: vec2<f32>,\n\
-         let i = gid.x;\n\
+         step: vec2<f32>,\n\
-         if i >= arrayLength(&probe_out) {{ return; }}\n\
+         radius: f32,\n\
-         let p = probe_in[i];\n\
+         }}\n\
-         probe_out[i] = rounded_rect_coverage(p.xy, vec2({}, {}), vec2({}, {}), p.z);\n\
+         @group(0) @binding(0) var<uniform> probe: Probe;\n\
+         @vertex\n\
+         fn probe_vs(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {{\n\
+         var xy = array(vec2(-1.0, -1.0), vec2(3.0, -1.0), vec2(-1.0, 3.0));\n\
+         return vec4(xy[vi], 0.0, 1.0);\n\
+         }}\n\
+         @fragment\n\
+         fn probe_fs(@builtin(position) pos: vec4<f32>) -> @location(0) f32 {{\n\
+         let at = probe.origin + floor(pos.xy) * probe.step;\n\
+         return rounded_rect_coverage(at, probe.top_left, probe.bot_right, probe.radius);\n\
          }}\n",
         wgsl_fn("distance_from_rect"),
         wgsl_fn("rounded_rect_coverage"),
-        TOP_LEFT.x,
-        TOP_LEFT.y,
-        BOT_RIGHT.x,
-        BOT_RIGHT.y,
     )
 }
 

rigs/virtgpu-probe/virtgpu-probe.c

@@ -0,0 +1,133 @@
+// What this VM's virtio-gpu actually offers, asked of the kernel and the
+// driver rather than inferred from the host's qemu command line.
+//
+//     cc -O2 -o virtgpu-probe virtgpu-probe.c -I/usr/include/libdrm -ldrm -lvulkan
+//     ./virtgpu-probe
+//
+// Written 2026-09-08 for the question "Venus keeps causing problems, is
+// there something to do with qemu instead" (docs/RUST.md, "Venus went
+// away for an hour"). Three things, each of which was guessed wrong at
+// least once before being measured:
+//
+//   1. Which capsets the host offers. Capset 6 (DRM) is the "native
+//      context" one -- RADV running in the guest against a passed-through
+//      DRM context instead of Venus proxying every Vulkan call. Whether
+//      it is available is a host-side fact this is the only way to read
+//      from in here.
+//   2. Whether the Vulkan device has compute. It does, and assuming it
+//      did not sent one investigation down the wrong path: the "no
+//      compute" finding on record is about the *Android emulator's*
+//      SwiftShader GL path, a different machine entirely.
+//   3. Whether plain Vulkan device teardown crashes on this adapter. It
+//      does not -- which is what makes wgpu's teardown SIGSEGV wgpu's and
+//      not the driver's, and is the kind of claim that is worthless
+//      without the negative half.
+//
+// C rather than a Rust crate on purpose: two of the three are ioctl and
+// loader questions that a wgpu-shaped rig cannot ask, and this needs to
+// keep working when the thing under suspicion is wgpu itself.
+#include <errno.h>
+#include <fcntl.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <unistd.h>
+
+#include <drm/virtgpu_drm.h>
+#include <vulkan/vulkan.h>
+#include <xf86drm.h>
+
+static const char *capset_name(int id) {
+    switch (id) {
+    case 1: return "VIRGL";
+    case 2: return "VIRGL2";
+    case 3: return "GFXSTREAM_VULKAN";
+    case 4: return "VENUS";
+    case 5: return "CROSS_DOMAIN";
+    case 6: return "DRM (native context)";
+    default: return "unknown";
+    }
+}
+
+static int capsets(void) {
+    int fd = open("/dev/dri/renderD128", O_RDWR);
+    if (fd < 0) {
+        printf("capsets: cannot open /dev/dri/renderD128: %s\n", strerror(errno));
+        return 1;
+    }
+    struct drm_virtgpu_getparam gp;
+    uint64_t mask = 0;
+    memset(&gp, 0, sizeof gp);
+    gp.param = VIRTGPU_PARAM_SUPPORTED_CAPSET_IDs;
+    gp.value = (uint64_t)(uintptr_t)&mask;
+    int rc = drmIoctl(fd, DRM_IOCTL_VIRTGPU_GETPARAM, &gp);
+    close(fd);
+    if (rc) {
+        // Not a virtio-gpu at all, or a kernel without the param: say
+        // which, rather than printing an empty list that reads like "the
+        // host offers nothing".
+        printf("capsets: SUPPORTED_CAPSET_IDs unavailable: %s\n", strerror(errno));
+        return 1;
+    }
+    printf("capsets: bitmask 0x%llx\n", (unsigned long long)mask);
+    for (int i = 1; i <= 8; i++)
+        if (mask & (1ull << i)) printf("  %d: %s\n", i, capset_name(i));
+    if (!(mask & (1ull << 6)))
+        printf("  (no capset 6: native context needs host-side virglrenderer + qemu support)\n");
+    return 0;
+}
+
+static int devices(void) {
+    VkInstanceCreateInfo ici = {.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO};
+    VkInstance inst;
+    if (vkCreateInstance(&ici, NULL, &inst) != VK_SUCCESS) {
+        printf("vulkan: no instance -- the loader found no usable ICD\n");
+        return 1;
+    }
+    uint32_t n = 0;
+    vkEnumeratePhysicalDevices(inst, &n, NULL);
+    if (n == 0) {
+        // The exact state this VM was in for an hour on 2026-09-08.
+        printf("vulkan: instance ok but ZERO devices -- the host refused a context\n");
+        vkDestroyInstance(inst, NULL);
+        return 1;
+    }
+    VkPhysicalDevice pd[8];
+    if (n > 8) n = 8;
+    vkEnumeratePhysicalDevices(inst, &n, pd);
+    for (uint32_t i = 0; i < n; i++) {
+        VkPhysicalDeviceProperties p;
+        vkGetPhysicalDeviceProperties(pd[i], &p);
+        printf("vulkan: %s (api %u.%u.%u)\n", p.deviceName, VK_VERSION_MAJOR(p.apiVersion),
+               VK_VERSION_MINOR(p.apiVersion), VK_VERSION_PATCH(p.apiVersion));
+        printf("  compute: %u invocations/workgroup, size %u,%u,%u, %u bytes shared\n",
+               p.limits.maxComputeWorkGroupInvocations, p.limits.maxComputeWorkGroupSize[0],
+               p.limits.maxComputeWorkGroupSize[1], p.limits.maxComputeWorkGroupSize[2],
+               p.limits.maxComputeSharedMemorySize);
+    }
+
+    // The negative half of "wgpu's teardown crashes on Venus": five
+    // devices and the instance, created and destroyed the plain way.
+    float prio = 1.0f;
+    VkDeviceQueueCreateInfo q = {.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO,
+                                 .queueFamilyIndex = 0, .queueCount = 1, .pQueuePriorities = &prio};
+    VkDeviceCreateInfo dci = {.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO,
+                              .queueCreateInfoCount = 1, .pQueueCreateInfos = &q};
+    for (int i = 0; i < 5; i++) {
+        VkDevice dev;
+        if (vkCreateDevice(pd[0], &dci, NULL, &dev) != VK_SUCCESS) {
+            printf("teardown: device %d could not be created\n", i);
+            return 1;
+        }
+        vkDestroyDevice(dev, NULL);
+    }
+    vkDestroyInstance(inst, NULL);
+    printf("teardown: 5 devices + instance created and destroyed cleanly\n");
+    return 0;
+}
+
+int main(void) {
+    int bad = capsets();
+    bad |= devices();
+    return bad;
+}