iris: stop requesting compute-shader limits nothing uses

adapter.request_device asked for Limits::default(), which requests
desktop-tier compute-shader limits unconditionally even though nothing in
iris/iris-core creates a ComputePipeline or writes a @compute stage. That
crashed device creation outright on a downlevel GL adapter reporting
OpenGL ES 3.0 (no compute at all) -- the Android emulator's
EMU_GPU=software/force-gles path, and any real GLES-3.0-only device.

New iris_core::device_limits(), shared by both platform backends, zeros
exactly the six max_compute_* fields rather than switching to a downlevel
Limits preset -- downlevel_webgl2_defaults() also zeros
max_storage_buffers_per_shader_stage, which shader.wgsl's vertex stage
needs. rigs/gpu-probe's own mirrored limits were updated to match.

Not verified against the actual SwiftShader-ES-3.0 crash on-device this
pass: the cold boot needed would have force-restarted this checkout's
emulator while another session had its own app running on it.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>

docs/DECISIONS.md

@@ -7,6 +7,48 @@ marked **DEFERRED** are ones the agent chose not to decide alone.
 
 ## 2026-09-05
 
+- **iris no longer asks every device for compute-shader limits it never
+  uses.** `adapter.request_device` (both `iris/src/android/render.rs` and
+  `iris/src/default/render.rs`) used `Limits::default()` plus an override
+  for `max_buffer_size`, and `Limits::default()` unconditionally requests
+  desktop-tier compute limits (`max_compute_workgroups_per_dimension:
+  65535`, per `wgpu_types`) even though nothing in `iris`/`iris-core`
+  creates a `ComputePipeline` or writes a `@compute` shader stage —
+  confirmed by grepping the whole tree, not assumed. That crashed
+  `request_device` outright on the Android emulator's software GL path
+  (`EMU_GPU=software`, `--features force-gles`): SwiftShader's GL reports
+  itself as OpenGL ES 3.0, which has no compute shaders at all, so the
+  adapter's real limit is 0 against the unconditional request for 65535 —
+  `RUST.md`'s "Software mode ... crashes for a third, different reason,"
+  2026-09-05, earlier today. The same would happen on any real
+  GLES-3.0-only Android device, not just the emulator. Fixed by a new
+  `iris_core::device_limits()` (`iris/core/src/render/mod.rs`), shared by
+  both platform backends so the two requests cannot drift, that zeros the
+  six `max_compute_*` fields explicitly rather than switching to a
+  downlevel `Limits` preset — `Limits::downlevel_webgl2_defaults()` was
+  considered and rejected: it also zeros
+  `max_storage_buffers_per_shader_stage`, and `shader.wgsl`'s vertex stage
+  reads four `var<storage>` buffers (rects, glyphs, masks, move_offsets),
+  so that preset would trade the compute crash for a bind-group-layout
+  one on the same downlevel hardware this is meant to support. No
+  capability check or fallback path was needed since nothing is being
+  disabled — the request is simply narrowed to what the pipeline actually
+  uses. `rigs/gpu-probe`'s own mirrored limits (it is deliberately its own
+  crate, not a workspace member, so it cannot call `device_limits()`
+  directly) were updated to match, and confirm `IRIS DEVICE: ok` against
+  this VM's own Vulkan and GL adapters. **Not verified this pass**: the
+  specific SwiftShader-ES-3.0 crash this fixes, on-device — the
+  `EMU_GPU=software` cold boot this needs would have force-restarted this
+  checkout's emulator while another session was actively running its own
+  app on it (`com.example.aiapp` had window focus at the time), so it was
+  left for a pass when the emulator is free rather than disrupting that
+  session. Everything reachable without the emulator is clean: `cargo
+  fmt`/`clippy --workspace --all-targets`/`test --workspace`, `cargo ndk
+  build`/`clippy` for `iris-android-app` with `force-gles`, and
+  `gpu-probe` against this VM's own Vulkan and GL(ES 3.2, which still has
+  compute and so would not have reproduced the crash even before this
+  fix — not a substitute for the real ES-3.0 test).
+
 - **P0's Compose half is built and smoke-tested on the emulator** — the
   `bench` build type, the shared `app/bench-fixture/` transcript, and an
   in-process fake backend (`BenchFixture.kt`/`BenchNetwork.kt`) that

docs/IRIS.md

@@ -8,6 +8,31 @@ capability that moved. Small and trivial changes do not go here.
 An entry gives the date, what changed, why, and a short before/after where
 it helps judge the change without the session that made it. Newest first.
 
+## 2026-09-05 (later still): `iris_core::device_limits()`, and iris no longer requests compute-shader limits
+
+New public function, `iris_core::device_limits() -> wgpu::Limits`. Why:
+`adapter.request_device`'s `required_limits` was `Limits::default()` plus
+a `max_buffer_size` override in both platform backends, and
+`Limits::default()` requests desktop-tier compute-shader limits
+unconditionally (`max_compute_workgroups_per_dimension: 65535`) even
+though nothing in `iris`/`iris-core` uses a `ComputePipeline` — that
+crashed device creation outright on a downlevel GL adapter reporting
+OpenGL ES 3.0 (no compute shaders at all: the Android emulator's
+`EMU_GPU=software` path, and any real GLES-3.0-only Android device).
+`device_limits()` is what both `android::render::AndroidRenderer::new`
+and `default::render::UiRenderer::new` now build their `required_limits`
+from, so the request cannot drift between the two backends.
+
+Before: `Limits { max_buffer_size: 1 << 30, ..Default::default() }`
+inlined in each backend. After: `iris_core::device_limits()`, which is
+the same thing with the six `max_compute_*` fields additionally zeroed.
+A caller building its own `DeviceDescriptor` outside these two backends
+(there are none today, but a third platform backend would want this)
+should call `device_limits()` rather than reaching for
+`Limits::default()` directly, unless it genuinely adds a compute pass —
+in which case it wants the specific compute limits that pass needs, not
+the desktop-tier default for everything.
+
 ## 2026-09-05 (later the same day): `iris_core::FrameReport` (RUST.md's I5 box)
 
 New public type, `iris_core::FrameReport` (re-exported from `iris_core`'s

docs/IRIS_TODO.md

@@ -7,6 +7,29 @@ order and what "done" looks like. Tick and date them in place.
 
 ## Fix
 
+- [x] **`request_device` asked for compute-shader limits it never uses
+  (2026-09-05).** `Limits::default()` (both `iris/src/android/render.rs`
+  and `iris/src/default/render.rs`) requests desktop-tier compute limits
+  unconditionally, even though nothing in `iris`/`iris-core` creates a
+  `ComputePipeline` or writes a `@compute` shader stage — confirmed by
+  grepping the whole tree, not assumed. That crashed device creation
+  outright on the Android emulator's software GL path (`EMU_GPU=software`,
+  `--features force-gles`): SwiftShader's GL reports itself as OpenGL ES
+  3.0, which has no compute shaders, so the adapter's real limit is 0
+  against the unconditional request for 65535 — the same would happen on
+  any real GLES-3.0-only Android device. Fixed by a new, shared
+  `iris_core::device_limits()` (`iris/core/src/render/mod.rs`) that zeros
+  exactly the six `max_compute_*` fields rather than switching to a
+  downlevel `Limits` preset — `downlevel_webgl2_defaults()` also zeros
+  `max_storage_buffers_per_shader_stage`, which `shader.wgsl`'s vertex
+  stage needs (four `var<storage>` buffers), so that preset would trade
+  this crash for a bind-group-layout one on the same hardware.
+  `rigs/gpu-probe`'s own hand-mirrored `Limits` (it is deliberately its
+  own crate, not able to call `device_limits()` directly) was updated to
+  match. See `DECISIONS.md` and RUST.md's I5 box for the account,
+  including what could not be re-verified on-device this pass (the
+  emulator was in concurrent use by another session).
+
 - [x] **Input does not fall through by input type (2026-09-04).**
   `SensorUi::run_sensors` (`src/default/sense.rs`) used to set "consumed,
   stop checking lower layers" from mere hover — a widget registered for

docs/RUST.md

@@ -59,6 +59,21 @@ session spending an afternoon on them again.
   end-to-end this pass. The fix itself is verified by direct, targeted
   logcat traces taken before that interference began, not by the
   aggregate script.
+- **iris no longer requests compute-shader limits it never uses,
+  2026-09-05.** `adapter.request_device`'s `Limits::default()` asks for
+  desktop-tier compute limits unconditionally even though nothing in
+  `iris`/`iris-core` uses a `ComputePipeline` -- confirmed by grep, not
+  assumed -- which is what crashed `request_device` outright under
+  `EMU_GPU=software`'s `force-gles` path (SwiftShader's GL reports OpenGL
+  ES 3.0, no compute at all). New shared `iris_core::device_limits()`
+  zeros exactly the six compute fields; `rigs/gpu-probe`'s own mirrored
+  limits were updated and confirm `IRIS DEVICE: ok` on this VM's own
+  Vulkan and GL adapters. **The specific SwiftShader-ES-3.0 crash this
+  fixes was not re-verified on-device this pass** -- the cold boot needed
+  would have force-restarted this checkout's emulator while another
+  session had its own app focused on it, so it was left rather than
+  disrupted. See this box's "Fixed, 2026-09-05, later the same day"
+  subsection (under the software-mode crash it fixes) and `DECISIONS.md`.
 - **Decided 2026-09-05: iris over Masonry**, by Iris, from the host-GPU
   numbers in I5's box and E1/E2's findings. See the Recommendation's item
   3 and `DECISIONS.md`. Next: the remaining screens and the app on iris —
@@ -3211,6 +3226,45 @@ silently on real hardware.
       general) explains the ~80-150ms software-mode numbers, since no GLES
       number under software mode could be taken at all.
 
+      **Fixed, 2026-09-05, later the same day.** Not "requesting compute
+      limits only when the adapter reports them" (a capability check with
+      a fallback) -- simpler than that, because iris has no code path that
+      needs compute at all: grepped the whole `iris`/`iris-core` tree for
+      `ComputePipeline`/`@compute` and found none, so the right fix is to
+      stop asking for compute limits, full stop, rather than to build a
+      fallback for a capability nothing uses. `iris_core::device_limits()`
+      (`iris/core/src/render/mod.rs`) is the one place both platform
+      backends now build their `required_limits` from: `Limits::default()`
+      with the six `max_compute_*` fields zeroed and `max_buffer_size`
+      still raised, as before. `Limits::downlevel_webgl2_defaults()` was
+      the first thing tried and rejected -- it also zeros
+      `max_storage_buffers_per_shader_stage`, and `shader.wgsl`'s vertex
+      stage reads four `var<storage>` buffers, so it would have traded
+      this crash for a bind-group-layout one on the same hardware.
+      `rigs/gpu-probe`'s own `Limits` (necessarily a hand-mirrored copy --
+      that rig is deliberately its own crate, not a workspace member) was
+      updated to match and re-run: `IRIS DEVICE: ok` against this VM's own
+      Vulkan (Venus) and GL (virgl, reports OpenGL ES 3.2) adapters.
+      **Not verified against the actual SwiftShader-ES-3.0 failure this
+      pass**: the `EMU_GPU=software` cold boot needed to reproduce it would
+      have force-restarted this checkout's shared emulator while another
+      session had `com.example.aiapp` focused and running on it (`adb
+      shell dumpsys window`), so this pass left that measurement rather
+      than disrupting concurrent work -- matching AGENTS.md's "coordinate
+      with peer agents" guidance rather than contending for the emulator.
+      Everything else: `cargo fmt --all`/`clippy --workspace --all-targets`/
+      `test --workspace` clean, `cargo ndk build`/`clippy` for
+      `iris-android-app --features transcript-screen,force-gles` clean
+      (only the pre-existing unused-`tabs-ui`-dependency warning, unrelated
+      to this change). This also means the software-mode question two
+      boxes up is still open, for the same original reason plus this new
+      one: a GLES number under `EMU_GPU=software` still has not been
+      taken, now blocked on emulator availability rather than on the
+      crash. A future pass should cold-boot `EMU_GPU=software` once the
+      emulator is free, confirm `dev.iris.android.demo` no longer aborts
+      on `request_device`, and take the `iris-scroll.sh` FrameReport row
+      that pairs with this box's host-GPU one.
+
       **Verification, this update.** `cargo fmt --all` (no diff),
       `cargo clippy --workspace --all-targets` (no warnings from the new
       code; pre-existing `wgpu`/`winit`/`naga` future-incompat notices

iris/core/src/render/mod.rs

@@ -23,6 +23,48 @@ pub use primitive::*;
 
 const SHAPE_SHADER: &str = include_str!("./shader.wgsl");
 
+/// The `wgpu::Limits` both platform backends (`android::render::
+/// AndroidRenderer::new`, `default::render::UiRenderer::new`) ask
+/// `Adapter::request_device` for -- shared so the two copies cannot drift,
+/// per AGENTS.md's "write the logic once."
+///
+/// Built from `Limits::default()`, **not** a downlevel variant: the shader
+/// (`shader.wgsl`) reads four `var<storage>` buffers (rects, glyphs, masks,
+/// move_offsets) from the vertex stage, and `Limits::downlevel_webgl2_defaults()`
+/// zeroes `max_storage_buffers_per_shader_stage` along with the compute
+/// limits below -- switching to it would trade one `request_device` crash
+/// for a bind-group-layout one on the same downlevel hardware this is meant
+/// to support. `max_buffer_size` is raised for the growing instance/atlas
+/// buffers (`ArrBuf`, `GpuTextures`); everything else is `default()`'s
+/// desktop-tier value, unchanged.
+///
+/// The six `max_compute_*` fields are zeroed because nothing in this crate
+/// creates a `ComputePipeline` or writes a `@compute` shader stage --
+/// grepped for both across `iris`/`iris-core` before writing this, found
+/// none. `Limits::default()` requests desktop-tier compute limits
+/// unconditionally (`max_compute_workgroups_per_dimension: 65535`) even
+/// though nothing asks a device to actually support compute, which is what
+/// crashed `request_device` on the Android emulator's software GL path
+/// (`EMU_GPU=software`, `force-gles`): SwiftShader's GL reports itself as
+/// OpenGL ES 3.0, which has no compute shaders, so the adapter's real limit
+/// is 0 and the unconditional request fails outright
+/// (`RUST.md`'s "Software mode ... crashes for a third, different reason").
+/// The same would happen on a real GLES-3.0-only Android device. If a
+/// future change adds a compute pass, request the specific limits it needs
+/// here rather than reverting to the desktop-tier default for everything.
+pub fn device_limits() -> Limits {
+    Limits {
+        max_buffer_size: 1 << 30,
+        max_compute_workgroup_storage_size: 0,
+        max_compute_invocations_per_workgroup: 0,
+        max_compute_workgroup_size_x: 0,
+        max_compute_workgroup_size_y: 0,
+        max_compute_workgroup_size_z: 0,
+        max_compute_workgroups_per_dimension: 0,
+        ..Default::default()
+    }
+}
+
 pub struct UiRenderNode {
     uniform_group: BindGroup,
     primitive_layout: BindGroupLayout,

iris/src/android/render.rs

@@ -86,12 +86,11 @@ impl AndroidRenderer {
 
         // Same request as the winit backend's `UiRenderer::new` -- no
         // binding-array features, see TEXTURES.md's "Recommended shape".
+        // `iris_core::device_limits()` is shared between the two backends;
+        // see its own doc for why it is not simply `Limits::default()`.
         let (device, queue) = adapter
             .request_device(&DeviceDescriptor {
-                required_limits: Limits {
+                required_limits: iris_core::device_limits(),
-                    max_buffer_size: 1 << 30,
-                    ..Default::default()
-                },
                 ..Default::default()
             })
             .block_on()

iris/src/default/render.rs

@@ -102,13 +102,12 @@ impl UiRenderer {
         // needs descriptor indexing. See TEXTURES.md's "Recommended shape"
         // for why the old binding array asked for
         // VK_EXT_descriptor_indexing unconditionally and did not survive a
-        // real share of Android GPUs.
+        // real share of Android GPUs. `iris_core::device_limits()` is
+        // shared with the Android backend; see its own doc for why it is
+        // not simply `Limits::default()`.
         let (device, queue) = adapter
             .request_device(&DeviceDescriptor {
-                required_limits: Limits {
+                required_limits: iris_core::device_limits(),
-                    max_buffer_size: 1 << 30,
-                    ..Default::default()
-                },
                 ..Default::default()
             })
             .block_on()

rigs/gpu-probe/src/main.rs

@@ -35,6 +35,34 @@ fn iris_features() -> Features {
 /// kept for the big storage buffers behind rects/glyphs.
 const IRIS_MAX_BUFFER_SIZE: u64 = 1 << 30;
 
+/// Mirrors `iris_core::device_limits()` (`iris/core/src/render/mod.rs`) --
+/// cannot call it directly, since this rig is deliberately its own crate,
+/// not a workspace member (this file's own Cargo.toml comment). Keep the
+/// two in sync by hand when one changes; this rig's whole purpose is "does
+/// the device iris actually builds come back," so a stale copy here would
+/// silently stop answering that question. Zeroed rather than left at
+/// `Limits::default()`'s desktop-tier values because nothing in iris
+/// creates a `ComputePipeline` or a `@compute` shader stage -- found by
+/// grepping the whole `iris`/`iris-core` tree before this rig's comment was
+/// written -- and the unconditional default request is what crashed
+/// `request_device` on the Android emulator's software GL path
+/// (`EMU_GPU=software`, `force-gles`: SwiftShader's GL reports itself as
+/// OpenGL ES 3.0, which has no compute shaders at all, so the adapter's
+/// real limit is 0). The same would happen on a real GLES-3.0-only Android
+/// device.
+fn iris_limits() -> Limits {
+    Limits {
+        max_buffer_size: IRIS_MAX_BUFFER_SIZE,
+        max_compute_workgroup_storage_size: 0,
+        max_compute_invocations_per_workgroup: 0,
+        max_compute_workgroup_size_x: 0,
+        max_compute_workgroup_size_y: 0,
+        max_compute_workgroup_size_z: 0,
+        max_compute_workgroups_per_dimension: 0,
+        ..Default::default()
+    }
+}
+
 fn main() {
     vk::report();
 
@@ -104,10 +132,7 @@ fn main() {
     // limits requested, this is expected to succeed everywhere -- this rig
     // is what turned that from an assumption into a measurement, first on
     // this emulator's software Vulkan.
-    let wanted = Limits {
+    let wanted = iris_limits();
-        max_buffer_size: IRIS_MAX_BUFFER_SIZE,
-        ..Default::default()
-    };
     match pollster::block_on(adapter.request_device(&DeviceDescriptor {
         required_features: iris_features(),
         required_limits: wanted,