iris: a mask is a shape, not a rectangle -- .masked_by, and touch obeys it

Iris, on the code fence: "the code block scrolling currently masks in an
inner rectangle. Ideally masks should have a shape associated with them,
rounded rectangle being one of them ... so that the mask becomes the
parent container with rounded edges. Make sure alpha works properly with
it, eg. on the corners where alpha should be decreased / multiplied."

`Mask` is now `{ primitive, parent }` -- the slot of a primitive already
written, plus the mask this one nests inside. The fragment stage
evaluates that primitive's own coverage at the masked pixel, through the
same `rounded_rect_coverage` a drawn rect goes through, and multiplies it
into the alpha along the whole `parent` chain. Nothing about the shape is
copied, so a rounded container's corner and its children's clipped corner
are one piece of arithmetic and cannot drift; two nested feathers dim a
pixel twice, which is the multiply she asked for.

`.masked()` is unchanged for callers: it writes an undrawn rect
(`Drawn::No`/`NOT_DRAWN` -- owned, moved, resized and freed like any
other primitive, simply never rasterized) and points at that, so square
clipping is the same mechanism rather than a special case. New
`.masked_by(shape)` draws `shape` behind the content in its own layer and
clips to the first primitive it drew, with no radius written twice; it
replaces `.masked().background(w)`, which drew both and clipped to the
box. `transcript-ui`'s `BlockFrame::Verbatim` is the first caller.

Hit-testing applies the shape (`SensorUi::run_sensors` ->
`UiRenderState::mask_admits`, coverage above one half, which is where the
drawn edge is), as well as the widget's own box -- the two ask different
questions and both have to hold. `primitive_corners` is a floor-for-floor
transliteration of the shader's `corners_of`, not `region.to_px()`: the
phone's 2.55 density puts nothing on a whole pixel, and skipping the
rounding disagrees with the pixels by up to one along each edge.

A mask's shape must be a rect, asserted by name in `set_mask_to`. A glyph
would need a CPU-side alpha plane before the hit test could agree with
the shader, and a standalone image a bind-group switch the fragment stage
cannot make. So no texture mask exists; the branch where one would go is
in both copies of `mask_coverage`. docs/LAYOUT.md's section end lists this
and the three other places the code is narrower than the design.

Tests. Layer 1, `layout_tests.rs`: the child's coverage swept across the
container's corner arc equals the container's own exactly; nested masks
multiply rather than intersect, asserted where both feathers are partial,
which is the only place the two differ; a press in a rounded-away corner
misses while one inside the curve and one on a straight edge hit; and
`a_plain_mask_still_clips_to_a_square_box`, the half this had no reason to
touch. The first version of the corner test swept the straight chord
between the arc's ends, which lies inside the circle everywhere -- it
proved nothing and said so, which is why it counts both sides now.

`iris/tests/mask_sdf.rs` is the only test here that needs a GPU: it lifts
`distance_from_rect` and `rounded_rect_coverage` out of
`iris_core::SHAPE_SHADER` by name -- lifted, not copied, since a copy
would be edited alongside the shader -- and runs them in a compute pass
over ~200k points at five radii against `iris_core::rounded_rect_coverage`.
Worst disagreement under 1e-5; the negative control (`+ 0.01` inside the
shader's smoothstep) fails it at 0.03.

Layer 2 for looking: `./run-headless.sh phone --phone --shot /tmp/mask.png
--seconds 6 -- -p transcript-fixture` draws the fixture's horizontally
scrolled code fence clipped on the curve at both top corners.

Two things found on the way and fixed here:

- The winit backend had the defect the Android one was fixed for in
  85869d0 -- `Backends::PRIMARY` and an `.expect` on the adapter. This
  VM's Venus device disappears when the host runs out of virgl contexts,
  which happened mid-task, and layer 2 aborted with `Could not get
  adapter!` while GL sat there working. It probes and rebuilds the
  instance on `Backends::GL` exactly as Android does now, and the request
  names the backends it tried. The rule had been written on one member of
  a set of two.
- `active_primitive_count` counted mask shapes, so `iris::frame`'s
  `primitives=` -- a number Iris reads off a phone report as "how much is
  on screen" -- would have gained one per masked widget.

`widget_trait!` now accepts a `///` doc comment on its functions, since
`masked_by` is public API and rustdoc is where a contract is read.

docs/LAYOUT.md, docs/RUST.md (both queue boxes, the commands, and where
the GPU test sits among the three layers), docs/IRIS.md and
docs/IRIS_TODO.md ("Masks defined relative to each other", now closed).

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

docs/IRIS.md

@@ -8,6 +8,48 @@ 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-08: masks have a shape -- `.masked_by(shape)`, and clipping applies to touch
+
+A mask no longer carries a rectangle. It carries **the slot of a
+primitive already drawn**, and the fragment stage evaluates that
+primitive's own coverage at each masked pixel and multiplies it into the
+alpha -- the same rounded-rect SDF the primitive itself is drawn with.
+Nothing about the shape is copied, so a rounded container's corner and
+the corner its content is cut to cannot fall out of step, and nested
+masks multiply rather than intersect: a pixel inside two feathered
+corners is dimmed by both.
+
+    // before -- the mask clipped to the padded box, the rounding was
+    // only painted behind it, and the two knew nothing of each other
+    field.scrollable_on(Axis::X)
+        .masked()
+        .pad(dp(FRAME_PAD_DP))
+        .background(rect(fill).radius(dp(FRAME_RADIUS_DP)))
+
+    // after -- one rect, drawn and clipped to
+    field.scrollable_on(Axis::X)
+        .pad(dp(FRAME_PAD_DP))
+        .masked_by(rect(fill).radius(dp(FRAME_RADIUS_DP)))
+
+`.masked()` is unchanged for callers and still clips to the widget's own
+box; under it, it now writes an undrawn rect primitive and points the
+mask at that, so square-cornered clipping is the same mechanism rather
+than a special case. `.masked_by(shape)` draws `shape` behind the
+content, in its own layer, and clips to the first primitive it drew.
+There is no radius or shape argument anywhere -- that is the point.
+
+**A press now has to be inside the shape, not just the box.** A corner
+the container rounded away is not there to be tapped, which needed the
+coverage function on the CPU as well as in the shader;
+`iris/tests/mask_sdf.rs` runs the shader's own text against the Rust one
+over a grid of points so the two cannot drift apart.
+
+One limit worth knowing before reaching for it: **a mask's shape must be
+a rect**, asserted by name. Clipping to a glyph or an image would need,
+respectively, a CPU-side alpha plane for the hit test and a bind-group
+switch the fragment stage cannot make. The shader has the branch where
+either would go.
+
 ## 2026-09-07: iris runs on a GLES-only Android device, and reports the renderer it cannot build
 
 `AndroidRenderer::new` asked wgpu for `Backends::PRIMARY`, which does not

docs/IRIS_TODO.md

@@ -701,12 +701,15 @@ Iris's report, verbatim, with a screenshot. Phone: Mali-G715 (Vulkan),
     has no rules for stays plain rather than being coloured by the
     nearest one's.
 
-- [ ] **Masks defined relative to each other.** Wanted: mask A multiplies
+- [x] **Masks defined relative to each other. (Done: chaining
-  by something *and also* applies mask B — a mask can reference a parent
+  2026-09-07 in d507ae4, the multiply 2026-09-08.)** Built exactly
-  mask, the way the move chain references a parent offset. Today masks
+  beside the move chain, as this asked: `Mask::parent` is a slot index
-  are independent regions. Design it beside the move chain (same shape:
+  and the fragment stage walks it under the same bound the move chain
-  a parent index and a bounded walk in the shader); do it when a real
+  uses. Each step multiplies the referenced primitive's coverage into
-  widget needs it, not before.
+  the pixel's alpha, so a pixel inside two feathered corners is dimmed
+  by both — the "multiplies by something *and also* applies mask B" half.
+  The real widget that needed it was the transcript's code fence inside
+  the list. See docs/LAYOUT.md's "Masks with a shape".
 - [ ] **Positions as a single float per scroll.** Iris raised, and half
   rejected, letting a scroll update one float rather than positions:
   input handling cares about most elements in a list, so absolute
@@ -1076,9 +1079,21 @@ do not duplicate it there.
 Iris pasted a full Copy report (Mali-G715 Vulkan, 2.55, 120Hz). What it
 showed, beyond her words:
 
-- [ ] **"Sometimes when I try to catch it while it's still moving
+- [x] **"Sometimes when I try to catch it while it's still moving
   (particularly if I drag) then it fails to stop & snap to where finger
-  is."** The report's release lines show catches ending as
+  is." (done 2026-09-07, b87f5a5.)** Built as described below.
+  `DragArbiter::press_start` takes a `PressState` -- what the target
+  looked like at the moment the press landed -- rather than asking the
+  list later, because by then the fling has already been cancelled and
+  the answer is no. The defect layer 1 found doing it: one touch-down
+  reaches every sensor under the finger, so a block and the tool row
+  containing it deliver the same `PressStart` twice, and re-reading the
+  state on the second delivery turned every catch back into an ordinary
+  slop-waiting press. Tests in
+  `iris/transcript-fixture/tests/catch_a_fling.rs`, with
+  `the_same_small_drag_on_a_settled_list_moves_nothing` as the half the
+  change had no reason to touch. **Not yet confirmed from the phone.**
+  The original reading follows. The report's release lines show catches ending as
   `v=-41`/`v=-274` pans, so the gesture *does* reach `Panning`, but the
   content under the finger does not follow it while the fling is still
   running and the slop has not been crossed. Compose: a down while

docs/LAYOUT.md

@@ -948,7 +948,7 @@ When this lands, copy this entry into `IRIS.md` (newest first):
 > design, the move-offset mechanism this shipped alongside, and the file
 > list.
 
-## Masks with a shape (decided 2026-09-07, not yet built)
+## Masks with a shape (decided 2026-09-07, built 2026-09-08)
 
 Iris, on the code block's scrolling: "the code block scrolling currently
 masks in an inner rectangle. Ideally masks should have a shape
@@ -1039,3 +1039,38 @@ and that the CPU SDF and the shader agree at a grid of points; a
 `run-headless.sh --phone` screenshot of a scrolled code block shows
 rounded corners with no square pixels poking out at the top and bottom
 of the scrolled content. Record the commands in RUST.md when it lands.
+
+### What was built (2026-09-08), and where it differs
+
+The commands and the screenshot are in docs/RUST.md's queue entry. Four
+places the code is narrower than the design above, each deliberate:
+
+- **No `kind` and no `flags` on `Mask`.** It is `{ primitive, parent }`.
+  The referenced instance already carries its own `binding`, so a copy
+  of it in the mask is a second thing to keep in step; *alpha only* is
+  the only mode there is, so there is nothing to select. Both are a
+  field away if a second mode appears.
+- **A mask's shape must be a rect.** `Painter::set_mask_to` asserts it,
+  by name, rather than leaving the shader to read a `rects` entry that
+  is not there. A glyph would need a CPU-side alpha plane before the
+  hit test could agree with the shader, and a standalone image needs a
+  bind-group switch the fragment stage cannot make (`masks_layout`'s own
+  comment on why an image's bind group must not name the masks buffer).
+  So **the texture-mask pass condition is not met and no texture mask
+  exists** — the point of the reference design is that adding one is a
+  binding check and a sampled alpha, with no new shader path, and the
+  shader's `mask_coverage` already has the branch where it would go.
+- **The shape is a primitive of its own, not always a drawn one.** A
+  plain `.masked()` writes an undrawn `RectPrimitive` at its region
+  (`Drawn::No`, `NOT_DRAWN`) and points the mask at that, so "clip to my
+  box" and "clip to that widget's rounded background" are one mechanism
+  and square-cornered clipping did not become a special case.
+  `.masked_by(shape)` draws `shape` behind the content — in its own
+  layer, the way `Stack` puts a background under its content — and
+  clips to the first primitive it drew.
+- **The CPU/shader agreement is a GPU test**, `iris/tests/mask_sdf.rs`,
+  the only test in the workspace that needs an adapter. It lifts
+  `distance_from_rect` and `rounded_rect_coverage` out of
+  `iris_core::SHAPE_SHADER` *by name* and runs them in a compute pass,
+  so the thing under test is the shader itself rather than a copy of it
+  that would be edited alongside.

docs/RUST.md

@@ -609,11 +609,12 @@ In order; two builders at a time. Each is ticked here by the agent that
 closes it.
 
 - [x] Test rig, layers 1 and 2 ("Three test layers" below), landed 2026-09-07.
-- [ ] Fling parity with Compose, and the phone's keyboard push-up, with
+- [x] Fling parity with Compose, and the phone's keyboard push-up, with
-  insets shown in the diagnostics overlay. The worktree note here was
+  insets shown in the diagnostics overlay. Every part is ticked in
-  stale by 2026-09-07 night: no worktree exists and the keyboard half
+  IRIS_TODO's 2026-09-07 entries: the keyboard half in the night entry,
-  is ticked in IRIS_TODO's night entry. What remains is the impulse
+  the estimator in the "later" one (Lsq2, not Impulse -- see the box
-  estimator item below.
+  below), and the catch in b87f5a5. **Not yet confirmed from the
+  phone**, which is what would close it for Iris rather than for us.
 - **Orchestrator note, 2026-09-07 late**: the tree was found holding a
   non-compiling diff from two killed agents (catch-a-fling in
   `sense.rs`/`selection.rs`; shaped masks in the render files). One
@@ -621,6 +622,11 @@ closes it.
   one sonnet agent owns report hygiene and the bench header in
   `bench_client.rs` and client-core's log ring. If both boxes below
   are still open and nothing is running, that work was cut off again.
+  **2026-09-08: it was cut off again** -- the sonnet agent's two boxes
+  had landed (7485d78, b8ea723), catch-a-fling had landed unticked
+  (b87f5a5), and the shaped-mask diff was left uncommitted in the tree
+  with one failing test. Both are closed below, by one session working
+  inline rather than by agents. Nothing is running now.
 - [x] Rows at the transcript's top edge: culled too early in one state,
   drawn through the header in the other (docs/IRIS_TODO.md, 2026-09-07).
   Done 2026-09-07, e922b73 + d507ae4; the root causes and the test names
@@ -648,9 +654,17 @@ closes it.
   `fonts.xml` monospace declaration against fontique's actually-scanned
   families, Android-only, verified `mono=Some("Droid Sans Mono")` on this
   checkout's emulator.
-- [ ] Scroll clamped at both ends, and Compose's impulse velocity
+- [x] Scroll clamped at both ends (e922b73, `List::clamp_to_content`)
-  estimator with min/max fling velocity (docs/IRIS_TODO.md, 2026-09-07
+  and Compose's velocity estimator (docs/IRIS_TODO.md, 2026-09-07
-  later). After the culling fix lands (same file).
+  later). Ticked 2026-09-08 against those entries, which were already
+  `[x]` while this box was not. Two things this box's own wording had
+  wrong, both corrected there by reading Compose's sources: the touch
+  path is **Lsq2 over absolute positions**, not `Strategy.Impulse`
+  (Impulse is the mouse-wheel/trackpad path), and there is **no minimum
+  fling velocity** on it -- `minimumFlingVelocity` belongs to
+  `NestedScrollInteropConnection`. So iris ports Lsq2, caps at 8000dp/s
+  and floors at 1px/s, and has no 50dp/s threshold Compose does not
+  have.
 - [x] **APK runtime logs in Dev Updater (Iris, 2026-09-07: "please add
   android / apk runtime log support to dev updater").** **Done
   2026-09-07** -- dev-updater `013d711`, and this repo's provider half;
@@ -716,9 +730,27 @@ closes it.
   APK installable (648 MB) -- RUST.md's logging section names both.
 - [ ] Input-event and timing instrumentation into the log ring, copied
   by the report button. After the logging route lands (same ring).
-- [ ] Catch-a-fling: down during a fling stops it at the down and drags
+- [x] **Catch-a-fling (done 2026-09-07, b87f5a5; verified and ticked
-  with no slop (docs/IRIS_TODO.md, night). Opus, next slot; uses the
+  2026-09-08).** A down on a moving list ends the fling on that sample
-  layer-1 harness.
+  and enters `Panning` with no `DRAG_SLOP` wait, which is Compose's
+  `scrollable(startDragImmediately = isScrollInProgress)`; a catch
+  released without moving is `Released(None)` rather than a `Tapped`,
+  which is also what Compose delivers. `DragArbiter::press_start` takes
+  a `PressState` -- what the target looked like when the press landed
+  (`already_selected`, `scrolling`) -- rather than reading the list
+  again later, because by then the fling it is asking about has already
+  been cancelled. The defect layer 1 found on the way: one touch-down
+  reaches every sensor under the finger, so a transcript row's block and
+  the tool row containing it deliver the same `PressStart` twice, and
+  re-reading `PressState` on the second delivery turned every catch back
+  into an ordinary slop-waiting press. Tests:
+  `transcript-fixture/tests/catch_a_fling.rs` (the recorded 120Hz flick,
+  150ms of fling, a down and three 2px moves -- the content tracks the
+  finger sample for sample, failing at the parent commit with "the
+  content 0.0px"), plus `the_same_small_drag_on_a_settled_list_moves_
+  nothing`, which is the half the change had no reason to touch: 6px is
+  inside `DRAG_SLOP`, so pinning the content on *every* press would pass
+  the first test and quietly take the slop away from every ordinary one.
 - [x] **Report hygiene (done 2026-09-07).** Ring takes Debug only from
   `iris`/`client_core` targets, Copy report always copies and trims the
   log. `client_core::log_ring::ring_accepts` is the one filter (Info+
@@ -806,15 +838,87 @@ closes it.
   renderer with the hook: `iris panic at .../render.rs:140:14: Could not
   get adapter!: NotFound {...}` in the ring on the run that died, and
   `iris app log: the previous run died -- ...` on the next one.
-- [ ] Masks with a shape -- docs/LAYOUT.md "Masks with a shape (decided
+- [x] **Masks with a shape (done 2026-09-08).** docs/LAYOUT.md's "Masks
-  2026-09-07)". A mask references a primitive already drawn
+  with a shape" carries the design and, at its end, the four places the
-  (rect SDF, texture or glyph alpha), chained and multiplied; `.masked()`
+  code is deliberately narrower than it. **Chaining landed early**,
-  points at the widget's own primitives; hit-testing applies the shape.
+  2026-09-07 (d507ae4). What landed now is the shape half and the
-  **Chaining landed early**, 2026-09-07 (d507ae4): a mask carries the
+  hit-testing half:
-  mask it was set inside and the fragment stage walks that chain, so
+
-  nesting works and `Painter::set_mask` no longer aborts on it. Still
+  `Mask` is `{ primitive, parent }` -- the slot of a primitive already
-  rectangles only -- the shape half, and the hit-testing half, are what
+  written, and the mask this one nests inside. The fragment stage
-  is left of this item.
+  evaluates that primitive's own coverage *at the masked pixel*
+  (`mask_coverage` in `shader.wgsl`, the same `rounded_rect_coverage` a
+  drawn rect goes through) and multiplies it into the pixel's alpha,
+  walking `parent` and multiplying every coverage on the chain. So the
+  container's corner and its children's clipped corner are one piece of
+  arithmetic, and two nested feathers dim a pixel twice -- the "alpha
+  should be decreased / multiplied" Iris asked for.
+
+  `.masked()` is unchanged for callers and writes an undrawn
+  `RectPrimitive` at its own region (`Drawn::No`/`NOT_DRAWN` -- owned,
+  moved, resized and freed like any other primitive, simply never
+  rasterized), so "clip to my box" and "clip to that rounded background"
+  are one mechanism rather than a square-cornered special case. New:
+  **`.masked_by(shape)`** draws `shape` behind the content in its own
+  layer and clips to the first primitive it drew, with no radius passed
+  twice -- it replaces `.masked().background(w)`, which drew the two and
+  clipped to the box. `transcript-ui`'s `BlockFrame::Verbatim` is its
+  first caller, which is the code fence Iris raised this about.
+
+  **Hit-testing applies the shape**: `SensorUi::run_sensors` asks
+  `UiRenderState::mask_admits` (coverage above one half, which is where
+  the drawn edge is) as well as the widget's own box -- the two ask
+  different questions and both have to hold. `primitive_corners` is a
+  `floor`-for-`floor` transliteration of the shader's `corners_of`,
+  which is the whole reason it is not `region.to_px()`: the phone's 2.55
+  density puts nothing on a whole pixel, and skipping the rounding
+  disagrees with the pixels by up to one along each edge.
+
+  A mask's shape **must be a rect**, asserted by name in
+  `Painter::set_mask_to`. A glyph would need a CPU-side alpha plane
+  before the hit test could agree with the shader, and a standalone
+  image a bind-group switch the fragment stage cannot make. So the
+  design's texture-mask pass condition is **not met and no texture mask
+  exists** -- the branch where one would go is in both `mask_coverage`s.
+
+  How it was checked, all four commands:
+
+      cd iris && cargo test --workspace          # layer 1 + the GPU test
+      cargo test -p iris --lib layout_tests::    # the four mask tests
+      cargo test -p iris --test mask_sdf         # CPU/shader agreement
+      ./run-headless.sh phone --phone --shot /tmp/mask.png --seconds 6 \
+          -- -p transcript-fixture               # layer 2, for looking
+
+  `iris/tests/mask_sdf.rs` is the only test in the workspace that needs
+  a GPU: it lifts `distance_from_rect` and `rounded_rect_coverage` out
+  of `iris_core::SHAPE_SHADER` **by name** and runs them in a compute
+  pass over a grid of ~200k points at five radii, against the CPU
+  `iris_core::rounded_rect_coverage` -- worst disagreement under 1e-5,
+  and the negative control (a `+ 0.01` inside the shader's smoothstep)
+  fails it at 0.03. It lifts rather than copies because a copy would be
+  edited alongside the shader, which is exactly the drift it exists to
+  catch. The layer-1 tests are in `layout_tests.rs`: the child's
+  coverage swept across the container's corner arc equals the
+  container's own *exactly* (the sweep goes from the arc's centre --
+  the straight chord between the arc's ends lies inside the circle
+  everywhere, so the first version of this test proved nothing and said
+  so); nested masks multiply rather than intersect, asserted where both
+  feathers are partial, which is the only place the two differ; a press
+  in a rounded-away corner misses while one inside the curve and one on
+  a straight edge hit; and `a_plain_mask_still_clips_to_a_square_box`,
+  the half the shape work had no reason to touch. The screenshot shows
+  the fixture's horizontally scrolled code fence clipped on the curve at
+  both top corners with no square pixels outside it.
+
+  **Also landed here, and not on this item's list**: the winit backend
+  had the same defect the Android one was fixed for in `85869d0` -- it
+  asked for a `Backends::PRIMARY` adapter and `.expect`ed one. This VM's
+  Venus device disappears whenever the host runs out of virgl contexts,
+  which happened mid-task, and layer 2 aborted with `Could not get
+  adapter!` while GL sat there working. `default::render::UiRenderer::
+  new` now probes and rebuilds the instance on `Backends::GL` exactly as
+  Android does, and the adapter request names the backends it tried. The
+  rule was written on one member of a set of two; this is the other.
 - [ ] Compose app: the `Reversed range` crash in `ToolInput.highlighted`
   (docs/TODO.md). Main branch, not rustify.
 
@@ -891,6 +995,19 @@ runs inside `cargo test`.
    layer 1 records that the platform was asked and layer 2 has no
    Android platform to ask.
 
+   **The one exception, added 2026-09-08**: `iris/tests/mask_sdf.rs`
+   needs a GPU but no compositor and no window -- it asks wgpu for an
+   adapter, runs two functions lifted out of `shader.wgsl` itself in a
+   compute pass, and compares the answers with the CPU transliteration
+   in `iris_core::render::sdf`. It sits inside `cargo test` because what
+   it checks is arithmetic rather than pixels: the fragment stage and
+   the hit test have to agree about where a rounded edge is, and neither
+   layer 1 (which cannot run the shader) nor layer 2 (where a
+   half-pixel disagreement is invisible) can say whether they do. Reach
+   for this shape only when the question is "do these two
+   implementations of one function agree" -- anything about what is
+   *drawn* is still layer 2.
+
 2. **A phone-shaped desktop window under headless sway -- for looking.**
 
        cd iris && ./run-headless.sh phone --phone --shot /tmp/p.png -- -p transcript-fixture

iris/Cargo.lock

@@ -1737,6 +1737,7 @@ dependencies = [
  "accesskit_winit",
  "android-view",
  "arboard",
+ "bytemuck",
  "image",
  "iris-core",
  "iris-macro",

iris/Cargo.toml

@@ -79,6 +79,9 @@ tokio = { workspace = true, features = ["sync", "rt", "rt-multi-thread", "time"]
 # package is fine -- cargo excludes dev-dependencies from the graph used
 # to build the library itself, so this only matters for `--examples`.
 tabs-ui = { path = "tabs-ui" }
+# `tests/mask_sdf.rs` only: the grid it hands the GPU and the coverages it
+# reads back. wgpu and pollster are ordinary dependencies already.
+bytemuck = { workspace = true }
 
 # Plain Instant-timed binaries, not criterion -- see benches/message_list.rs's
 # header for why. `harness = false` opts out of the unstable `#[bench]`

iris/core/src/render/data.rs

@@ -49,22 +49,47 @@ impl MaskIdx {
 
 pub type MoveIdx = Id<u32>;
 
+/// A clip, as a reference to a primitive already written plus the mask it
+/// nests inside. The fragment stage evaluates that primitive's coverage
+/// *at the masked pixel* -- for a rect, the same `rounded_rect_coverage`
+/// from the same SDF the rect itself is drawn with -- and multiplies it
+/// into the pixel's alpha, so a rounded container's corner and its
+/// children's clipped corner are the same arithmetic and cannot disagree.
+/// See LAYOUT.md's "Masks with a shape".
+///
+/// **No `kind` and no `flags`**, which the design sketched: the referenced
+/// instance already carries its own `binding`, and a copy of it here is a
+/// second thing to keep in step; alpha-only is the only mode there is, so
+/// there is nothing to select. Both are a field away if a second mode
+/// appears.
 #[repr(C)]
 #[derive(Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
 pub struct Mask {
-    pub region: UiRegion,
+    /// The slot in `UiRenderState::primitives` of the primitive whose
-    /// The mask-owning widget's own move slot -- resolved in the fragment
+    /// coverage this mask is. Today always a `RectPrimitive`: a glyph or
-    /// shader against the same chain the vertex shader walks for a
+    /// a standalone image would need, respectively, a CPU-side alpha
-    /// primitive's own corners, so a mask and the content clipped by it
+    /// plane for the hit test to agree with the shader, and a bind-group
-    /// can move independently. See LAYOUT.md section 2b.
+    /// switch the fragment stage cannot make -- `Painter::set_mask`
-    pub move_idx: MoveIdx,
+    /// rejects both by name rather than leaving the shader to read a rect
+    /// that is not there.
+    ///
+    /// Who owns it depends on which way the mask was set. A plain
+    /// `.masked()` writes its own undrawn rect, so the primitive is in
+    /// the masking widget's `ActiveData::primitives` and lives exactly as
+    /// long as the mask. `.masked_by(shape)` points at a *child's*
+    /// primitive, which that child can free on any redraw of its own --
+    /// so `UiRenderState::remask_shape_users` marks the mask's owner for
+    /// redraw whenever a referenced slot is freed, since that widget's
+    /// own `set_mask` is the only thing that resolves the slot again.
+    pub primitive: u32,
     /// The mask this one was set *inside* (`MaskIdx::NONE` at the top), so
-    /// clipping nests: the fragment stage walks the chain and a pixel has
+    /// clipping nests: the fragment stage walks the chain and multiplies
-    /// to be inside every mask on it. Chained rather than intersected on
+    /// every coverage on it, which is what makes a pixel inside two
-    /// the CPU because each mask moves with its own widget -- a code fence
+    /// feathered corners dimmed by both. Chained rather than intersected
-    /// inside a transcript row carries the row's scroll, the list's own
+    /// on the CPU because each mask moves with its own widget -- a code
-    /// box does not, and one region resolved when the fence was last drawn
+    /// fence inside a transcript row carries the row's scroll, the list's
-    /// gets the second of those wrong as soon as the row moves.
+    /// own box does not, and one region resolved when the fence was last
+    /// drawn gets the second of those wrong as soon as the row moves.
     ///
     /// A child holds one ref on its parent's slot (`Painter::set_mask`),
     /// released when the child's own slot goes

iris/core/src/render/mod.rs

@@ -18,6 +18,7 @@ mod atlas;
 mod data;
 mod frame_report;
 mod primitive;
+mod sdf;
 mod texture;
 mod util;
 
@@ -25,8 +26,14 @@ pub use atlas::*;
 pub use data::{Mask, MaskIdx, MoveIdx, MoveOffset};
 pub use frame_report::{FrameReport, FrameStats, JANK_THRESHOLD};
 pub use primitive::*;
+pub use sdf::{distance_from_rect, rounded_rect_coverage};
 
-const SHAPE_SHADER: &str = include_str!("./shader.wgsl");
+/// The one shader every primitive is drawn with. Public so a test can run
+/// a function out of it against the CPU transliteration in [`sdf`] --
+/// `iris/tests/mask_sdf.rs`, which LAYOUT.md's "Masks with a shape" turns
+/// on: a masked corner that cannot be tapped and a masked corner that is
+/// not drawn are only the same corner while the two agree.
+pub const SHAPE_SHADER: &str = include_str!("./shader.wgsl");
 
 /// The `wgpu::Limits` both platform backends (`android::render::
 /// AndroidRenderer::new`, `default::render::UiRenderer::new`) ask

iris/core/src/render/primitive.rs

@@ -21,6 +21,10 @@ pub const IMAGE_BINDING: u32 = 1;
 pub trait Primitive: Pod {
     const BINDING: u32;
     fn vec(data: &mut PrimitiveData) -> &mut PrimitiveVec<Self>;
+    /// The read-only half of [`Self::vec`], for a caller that wants to
+    /// look one entry up rather than write one -- a mask reading the
+    /// radius of the rect it clips to ([`Primitives::data`]).
+    fn vec_ref(data: &PrimitiveData) -> &PrimitiveVec<Self>;
 }
 
 macro_rules! primitives {
@@ -86,6 +90,9 @@ macro_rules! primitives {
                 fn vec(data: &mut PrimitiveData) -> &mut PrimitiveVec<Self> {
                     &mut data.$name
                 }
+                fn vec_ref(data: &PrimitiveData) -> &PrimitiveVec<Self> {
+                    &data.$name
+                }
             }
         )*
     };
@@ -287,6 +294,15 @@ impl Primitives {
         &self.instances[slot as usize]
     }
 
+    /// The per-primitive data behind `slot`, or `None` if that slot holds
+    /// a different kind of primitive -- the `binding` check is the same
+    /// one the shader's dispatch switch makes, and it is what stops a
+    /// caller reading a glyph's index into the rect table.
+    pub fn primitive_data<P: Primitive>(&self, slot: u32) -> Option<&P> {
+        let inst = self.instance(slot);
+        (inst.binding == P::BINDING).then(|| &P::vec_ref(&self.data)[inst.idx as usize])
+    }
+
     pub fn region_mut(&mut self, h: &PrimitiveHandle) -> &mut UiRegion {
         self.updated = true;
         &mut self.instances[h.slot as usize].region
@@ -391,10 +407,23 @@ pub struct OrderChange {
     pub pos: usize,
 }
 
+/// Whether a primitive goes into its layer's draw order. [`Drawn::No`] is
+/// a primitive written only to be *referenced* -- a mask's shape
+/// (LAYOUT.md's "Masks with a shape"). It is owned, moved, resized and
+/// freed exactly like any other; it is simply never rasterized.
+#[derive(Clone, Copy, PartialEq, Eq, Debug)]
+pub enum Drawn {
+    Yes,
+    No,
+}
+
+/// The `pos` of a [`Drawn::No`] primitive: it is in no layer's order, so
+/// there is no position to renumber or free.
+pub const NOT_DRAWN: usize = usize::MAX;
+
 /// Where one primitive lives: its stable slot in the global arena, and
-/// where in a layer's draw order it currently sits. A handle with no
+/// where in a layer's draw order it currently sits ([`NOT_DRAWN`] if it is
-/// layer position (`pos == NOT_DRAWN`) is a primitive that exists to be
+/// only referenced).
-/// *referenced* -- a mask's shape -- and is never rasterized.
 #[derive(Debug)]
 pub struct PrimitiveHandle {
     pub layer: usize,

iris/core/src/render/sdf.rs

@@ -0,0 +1,54 @@
+//! The rounded-rect coverage function, on the CPU.
+//!
+//! `shader.wgsl`'s `distance_from_rect`/`rounded_rect_coverage` are a
+//! transliteration of these two, line for line, and
+//! `mask_sdf_matches_the_shader` in `iris`'s layout tests compares the two
+//! at a grid of points against values the shader itself produced. They are
+//! kept together here, in the crate both a renderer and a hit test can
+//! reach, because LAYOUT.md's "Masks with a shape" turns on the two
+//! agreeing: a masked corner that cannot be tapped and a masked corner
+//! that is not drawn have to be the same corner, and they are only the
+//! same corner while one function decides both.
+//!
+//! Window pixels throughout, matching the shader's `pos` -- not `UiRegion`
+//! units, which the shader has already resolved by the time it evaluates
+//! this.
+
+use crate::util::Vec2;
+
+/// The signed distance from `pos` to a rounded rect given by its centre,
+/// its corner offset (half its size) and its corner `radius`. Negative
+/// inside.
+pub fn distance_from_rect(pos: Vec2, center: Vec2, corner: Vec2, radius: f32) -> f32 {
+    // vec from center to pixel
+    let p = pos - center;
+    // vec from inner rect corner to pixel
+    let q = Vec2::new(
+        p.x.abs() - (corner.x - radius),
+        p.y.abs() - (corner.y - radius),
+    );
+    let clamped = Vec2::new(q.x.max(0.0), q.y.max(0.0));
+    (clamped.x * clamped.x + clamped.y * clamped.y).sqrt() - radius
+}
+
+/// How much of the pixel at `pos` a rounded rect covers, anti-aliased over
+/// the half-pixel either side of its edge: 1 well inside, 0 well outside.
+///
+/// The half-pixel feather is why a hit test asks for **more than a half**
+/// rather than "any coverage at all": half is where the geometric edge is,
+/// so the two answer the same question the drawn shape does.
+pub fn rounded_rect_coverage(pos: Vec2, top_left: Vec2, bot_right: Vec2, radius: f32) -> f32 {
+    let edge: f32 = 0.5;
+    let corner = (bot_right - top_left) / 2.0;
+    let center = top_left + corner;
+    let dist = distance_from_rect(pos, center, corner, radius);
+    1.0 - smoothstep(-edge.min(radius), edge, dist)
+}
+
+/// WGSL's `smoothstep`, which Rust has no equivalent of. Undefined in WGSL
+/// when `low == high`, which is why the caller above never passes a zero
+/// radius into the low edge without `edge` bounding it.
+fn smoothstep(low: f32, high: f32, x: f32) -> f32 {
+    let t = ((x - low) / (high - low)).clamp(0.0, 1.0);
+    t * t * (3.0 - 2.0 * t)
+}

iris/core/src/render/shader.wgsl

@@ -30,13 +30,11 @@ struct GlyphInfo {
     flags: u32,
 }
 
+/// Mirrors `Mask` in data.rs: the slot of the primitive whose coverage
+/// clips this mask's subtree, and the mask it nests inside
+/// (`4294967295u` at the top).
 struct Mask {
-    x: UiSpan,
+    primitive: u32,
-    y: UiSpan,
-    move_idx: u32,
-    /// The mask this one is nested inside, or `4294967295u`. Mirrors
-    /// `Mask::parent` in data.rs; walked below with the same bound the
-    /// move chain uses.
     parent: u32,
 }
 
@@ -57,11 +55,6 @@ struct UiScalar {
     abs: f32,
 }
 
-struct UiVec2 {
-    rel: vec2<f32>,
-    abs: vec2<f32>,
-}
-
 // The shared glyph atlas: every page is one layer. Growing it recreates this
 // texture with headroom and copies the old layers across -- see
 // GpuTextures::grow_array -- rather than the binding_array<texture_2d<f32>>
@@ -157,6 +150,27 @@ struct Region {
     bot_right: vec2<f32>,
 }
 
+/// One primitive's on-screen corners in window pixels. Written once and
+/// used by both stages: the vertex stage for the primitive it is drawing,
+/// the fragment stage for a mask's -- so the shape a mask clips to and the
+/// shape that was drawn cannot be computed two different ways.
+struct Corners {
+    top_left: vec2<f32>,
+    bot_right: vec2<f32>,
+}
+
+fn corners_of(inst: PrimitiveInstance) -> Corners {
+    let top_left_rel = vec2(inst.x.start.rel, inst.y.start.rel);
+    let top_left_abs = vec2(inst.x.start.abs, inst.y.start.abs);
+    let bot_right_rel = vec2(inst.x.end.rel, inst.y.end.rel);
+    let bot_right_abs = vec2(inst.x.end.abs, inst.y.end.abs);
+    let move_delta = resolve_move(inst.move_idx);
+    return Corners(
+        floor(top_left_rel * window.dim) + floor(top_left_abs) + move_delta,
+        floor(bot_right_rel * window.dim) + floor(bot_right_abs) + move_delta,
+    );
+}
+
 @vertex
 fn vs_main(
     @builtin(vertex_index) vi: u32,
@@ -165,14 +179,9 @@ fn vs_main(
     var out: VertexOutput;
     let inst = instances[in.slot];
 
-    let top_left_rel = vec2(inst.x.start.rel, inst.y.start.rel);
+    let c = corners_of(inst);
-    let top_left_abs = vec2(inst.x.start.abs, inst.y.start.abs);
+    let top_left = c.top_left;
-    let bot_right_rel = vec2(inst.x.end.rel, inst.y.end.rel);
+    let bot_right = c.bot_right;
-    let bot_right_abs = vec2(inst.x.end.abs, inst.y.end.abs);
-
-    let move_delta = resolve_move(inst.move_idx);
-    let top_left = floor(top_left_rel * window.dim) + floor(top_left_abs) + move_delta;
-    let bot_right = floor(bot_right_rel * window.dim) + floor(bot_right_abs) + move_delta;
     let size = bot_right - top_left;
 
     let uv = vec2<f32>(
@@ -214,28 +223,38 @@ fn fs_main(
         }
     }
     // Every mask on the chain, not just the innermost: a widget that set
-    // its own mask inside another is clipped by both, and each carries its
+    // its own mask inside another is clipped by both, and the coverages
-    // own move slot (`Mask::parent` in data.rs).
+    // multiply -- so a pixel inside two feathered corners is dimmed by
+    // both, which is what a compositor does (`Mask::parent` in data.rs).
     var mask_idx = in.mask_idx;
     for (var step = 0u; step < PARENT_CHAIN_LIMIT; step++) {
         if mask_idx == 4294967295u {
             break;
         }
         let mask = masks[mask_idx];
-        let mask_delta = resolve_move(mask.move_idx);
+        color.a *= mask_coverage(pos, mask);
-        let tl = UiVec2(vec2(mask.x.start.rel, mask.y.start.rel), vec2(mask.x.start.abs, mask.y.start.abs));
-        let br = UiVec2(vec2(mask.x.end.rel, mask.y.end.rel), vec2(mask.x.end.abs, mask.y.end.abs));
-
-        let top_left = floor(tl.rel * window.dim) + floor(tl.abs) + mask_delta;
-        let bot_right = floor(br.rel * window.dim) + floor(br.abs) + mask_delta;
-        if pos.x < top_left.x || pos.x > bot_right.x || pos.y < top_left.y || pos.y > bot_right.y {
-            color *= 0.0;
-        }
         mask_idx = mask.parent;
     }
     return color;
 }
 
+/// How much of `pos` one mask lets through: the referenced primitive's
+/// own coverage at that pixel, from the same SDF the primitive is drawn
+/// with. Nothing about the shape is copied into the mask, so a rounded
+/// container's corner and its children's clipped corner are the same
+/// arithmetic.
+fn mask_coverage(pos: vec2<f32>, mask: Mask) -> f32 {
+    let inst = instances[mask.primitive];
+    if inst.binding != RECT {
+        // Unreachable: `Painter::set_mask` rejects a glyph or an image
+        // shape by name (see `Mask::primitive`). Letting the pixel
+        // through rather than reading a `rects` entry that is not there.
+        return 1.0;
+    }
+    let c = corners_of(inst);
+    return rounded_rect_coverage(pos, c.top_left, c.bot_right, rects[inst.idx].radius);
+}
+
 fn draw_texture(region: Region) -> vec4<f32> {
     return textureSample(image_texture, samp, region.uv);
 }
@@ -251,19 +270,35 @@ fn draw_glyph(region: Region, g: GlyphInfo) -> vec4<f32> {
     return color;
 }
 
+/// The anti-aliased coverage of a rounded rect at one pixel -- the one
+/// function both a drawn rect and a mask go through, and the
+/// transliteration of `iris_core::rounded_rect_coverage` on the CPU,
+/// which the hit test uses so a corner that cannot be tapped and a corner
+/// that is not drawn are the same corner.
+fn rounded_rect_coverage(
+    pos: vec2<f32>,
+    top_left: vec2<f32>,
+    bot_right: vec2<f32>,
+    radius: f32,
+) -> f32 {
+    let edge = 0.5;
+    let corner = (bot_right - top_left) / 2.0;
+    let center = top_left + corner;
+    let dist = distance_from_rect(pos, center, corner, radius);
+    return 1.0 - smoothstep(-min(edge, radius), edge, dist);
+}
+
 fn draw_rounded_rect(region: Region, rect: Rect) -> vec4<f32> {
     var color = unpack4x8unorm(rect.color);
 
     let edge = 0.5;
 
-    let size = region.bot_right - region.top_left;
+    color.a *= rounded_rect_coverage(region.pos, region.top_left, region.bot_right, rect.radius);
-    let corner = size / 2.0;
-    let center = region.top_left + corner;
-
-    let dist = distance_from_rect(region.pos, center, corner, rect.radius);
-    color.a *= 1.0 - smoothstep(-min(edge, rect.radius), edge, dist);
 
     if rect.thickness > 0.0 {
+        let size = region.bot_right - region.top_left;
+        let corner = size / 2.0;
+        let center = region.top_left + corner;
         let dist2 = distance_from_rect(region.pos, center, corner - rect.thickness, rect.inner_radius);
         color.a *= smoothstep(-min(edge, rect.inner_radius), edge, dist2);
     }

iris/core/src/ui/painter.rs

@@ -1,7 +1,10 @@
 use crate::{
-    RenderedText, Size, StrongWidget, TextAttrs, TextBuffer, TextData, TextureHandle, UiRegion,
+    Color, RenderedText, Size, StrongWidget, TextAttrs, TextBuffer, TextData, TextureHandle,
-    UiRenderState, UiRsc, UiScalar, UiVec2, WidgetId,
+    UiRegion, UiRenderState, UiRsc, UiScalar, UiVec2, WidgetId,
-    render::{GlyphPrimitive, Mask, MaskIdx, MoveIdx, Primitive, PrimitiveHandle, PrimitiveInst},
+    render::{
+        Drawn, GlyphPrimitive, Mask, MaskIdx, MoveIdx, Primitive, PrimitiveHandle, PrimitiveInst,
+        RectPrimitive,
+    },
     util::Vec2,
 };
 
@@ -26,8 +29,20 @@ pub struct Painter<'a> {
 
 impl<'a> Painter<'a> {
     fn primitive_at<P: Primitive>(&mut self, primitive: P, region: UiRegion) {
+        self.write_primitive(primitive, region, Drawn::Yes);
+    }
+
+    /// The one path every primitive this widget owns goes through --
+    /// drawn or, for a mask's shape, only referenced.
+    fn write_primitive<P: Primitive>(
+        &mut self,
+        primitive: P,
+        region: UiRegion,
+        drawn: Drawn,
+    ) -> u32 {
         let h = self.state.write_primitive(
             self.layer,
+            drawn,
             PrimitiveInst {
                 id: self.id,
                 primitive,
@@ -40,7 +55,9 @@ impl<'a> Painter<'a> {
             // TODO: I have no clue if this works at all :joy:
             self.rsc.ui_mut().masks.push_ref(self.mask);
         }
+        let slot = h.slot;
         self.primitives.push(h);
+        slot
     }
 
     /// Writes a primitive to be rendered
@@ -55,16 +72,46 @@ impl<'a> Painter<'a> {
     /// Clip everything this widget draws, itself and its descendants, to
     /// `region`. One call per widget; a widget drawn inside another
     /// widget's mask nests instead -- the new mask chains to the inherited
-    /// one (`Mask::parent`) and the fragment stage requires a pixel to be
+    /// one (`Mask::parent`) and the fragment stage multiplies both
-    /// inside both, which is what lets a transcript row's code fence clip
+    /// coverages, which is what lets a transcript row's code fence clip
     /// to itself *and* to the list it scrolls inside.
     ///
+    /// The clip is a **primitive**, not a rectangle copied into the mask:
+    /// this writes an undrawn `RectPrimitive` at `region` and points the
+    /// mask at it, so the fragment stage evaluates the same rounded-rect
+    /// coverage a drawn rect gets. See LAYOUT.md's "Masks with a shape".
+    ///
     /// The slot is allocated once and **rewritten in place** on every
     /// later draw rather than pushed again, because a descendant whose own
     /// region did not change is not redrawn (`draw_inner`'s fast path) and
     /// so keeps pointing at whichever slot it was drawn under. See
     /// `ActiveData::own_mask` for what pushing a fresh one cost.
     pub fn set_mask(&mut self, region: UiRegion) {
+        let shape = self.write_primitive(RectPrimitive::color(Color::NONE), region, Drawn::No);
+        self.set_mask_to(shape);
+    }
+
+    /// Clip everything this widget draws after this call to `shape`'s
+    /// own shape -- the first primitive `shape`'s subtree drew, which
+    /// must already have been drawn this frame
+    /// (`UiRenderState::first_primitive`). What `.masked_by()` uses to
+    /// clip a container's content to the rounded background it draws,
+    /// with no radius argument anywhere that could fall out of step with
+    /// the one being drawn.
+    pub fn set_mask_to_widget<W: ?Sized>(&mut self, shape: &StrongWidget<W>) {
+        let slot = self.state.first_primitive(shape.id()).unwrap_or_else(|| {
+            panic!(
+                "'{}' was given as a mask's shape but drew no primitive, so there is nothing to \
+                 clip to",
+                self.rsc.widgets().label(shape.id()),
+            )
+        });
+        self.set_mask_to(slot);
+    }
+
+    /// Points this widget's mask at a primitive that has already been
+    /// written -- the shared half of [`Self::set_mask`].
+    fn set_mask_to(&mut self, shape: u32) {
         // `assert!`, not `debug_assert!`: one comparison per widget draw,
         // and the second call silently *replacing* the first is a widget
         // drawn unclipped -- which reaches the screen and nothing says so.
@@ -75,10 +122,20 @@ impl<'a> Painter<'a> {
             "set_mask called twice while drawing one widget: the second would replace the first \
              rather than nest inside it",
         );
+        // A glyph would need a CPU-side alpha plane for the hit test to
+        // agree with the shader, and a standalone image a bind-group
+        // switch the fragment stage cannot make -- see `Mask::primitive`.
+        // Named here rather than left to the shader, which would read a
+        // rect that is not there and clip to nothing.
+        let binding = self.state.primitives.instance(shape).binding;
+        assert_eq!(
+            binding,
+            RectPrimitive::BINDING,
+            "a mask's shape must be a rect primitive; primitive {shape} is binding {binding}",
+        );
         let parent = self.mask;
         let mask = Mask {
-            region,
+            primitive: shape,
-            move_idx: self.move_slot,
             parent,
         };
         let old_parent = if self.own_mask == MaskIdx::NONE {

iris/core/src/ui/render_state.rs

@@ -4,7 +4,10 @@ use std::time::{Duration, Instant};
 use crate::{
     ActiveData, IdLike, MaskIdx, MoveIdx, Painter, PixelRegion, PrimitiveLayers, RegionAlign,
     StrongWidget, UiRegion, UiRsc, UiVec2, WidgetId, Widgets,
-    render::{MoveOffset, Primitive, PrimitiveHandle, PrimitiveInst, Primitives},
+    render::{
+        Drawn, MoveOffset, NOT_DRAWN, Primitive, PrimitiveHandle, PrimitiveInst, Primitives,
+        RectPrimitive, rounded_rect_coverage,
+    },
     util::{HashMap, HashSet, Id, Vec2},
 };
 
@@ -166,14 +169,19 @@ impl UiRenderState {
         )
     }
 
-    /// Writes a primitive into the arena and into `layer`'s draw order.
+    /// Writes a primitive into the arena and, unless it is
+    /// [`Drawn::No`], into `layer`'s draw order.
     pub(super) fn write_primitive<P: Primitive>(
         &mut self,
         layer: usize,
+        drawn: Drawn,
         inst: PrimitiveInst<P>,
     ) -> PrimitiveHandle {
         let (slot, data_idx) = self.primitives.alloc(inst);
-        let pos = self.layers[layer].push(slot, false);
+        let pos = match drawn {
+            Drawn::Yes => self.layers[layer].push(slot, false),
+            Drawn::No => NOT_DRAWN,
+        };
         PrimitiveHandle {
             layer,
             pos,
@@ -357,8 +365,16 @@ impl UiRenderState {
     /// on screen", which is what a report reads as "did this frame have
     /// more to draw than the last one", not "how much work did this frame
     /// do" (`take_counters` answers that).
+    ///
+    /// A mask's shape does not count: it is a [`Drawn::No`] primitive
+    /// that is never rasterized, so including it would put one extra on
+    /// the line for every masked widget and make a number Iris reads off
+    /// a phone report disagree with what is drawn.
     pub fn active_primitive_count(&self) -> usize {
-        self.active.values().map(|a| a.primitives.len()).sum()
+        self.active
+            .values()
+            .map(|a| a.primitives.iter().filter(|h| h.pos != NOT_DRAWN).count())
+            .sum()
     }
 
     fn redraw_all(&mut self, root: Option<&StrongWidget>, rsc: &mut dyn UiRsc) {
@@ -690,11 +706,14 @@ impl UiRenderState {
         if let Some(active) = &mut active {
             for h in &active.primitives {
                 let mask = self.primitives.free(h);
-                self.layers[h.layer].free(h.pos, h.is_image());
+                if h.pos != NOT_DRAWN {
+                    self.layers[h.layer].free(h.pos, h.is_image());
+                }
                 if mask != MaskIdx::NONE {
                     rsc.ui_mut().masks.remove(mask);
                 }
             }
+            Self::remask_shape_users(&self.active, id, active.own_mask, &active.primitives, rsc);
             active.textures.clear();
             rsc.ui_mut().textures.free();
             if undraw {
@@ -735,6 +754,54 @@ impl UiRenderState {
         active
     }
 
+    /// A mask whose shape primitive was just freed clips to a slot that
+    /// now holds something else, so the widget that owns it is marked for
+    /// redraw -- its own `set_mask` is the only thing that resolves the
+    /// slot, and it is the same mechanism a dirty widget already goes
+    /// through.
+    ///
+    /// `own` is the mask belonging to the widget being removed and is
+    /// skipped: this runs in the middle of that widget's own redraw,
+    /// which sets its mask again on the way out, and a mark left on
+    /// itself would redraw it every frame from then on. Skipping it is
+    /// also what keeps the O(active) scan off the ordinary path -- a
+    /// plain `.masked()` frees exactly its own shape, so `stale` is empty
+    /// and this returns before touching `active`.
+    ///
+    /// Both `Vec`s start empty and stay unallocated in that case, and
+    /// membership is a linear scan of two lists that are a handful long
+    /// (a widget's own primitives, and the live masks): this runs once
+    /// per widget removed, which is once per dirty widget per frame, and
+    /// a set built there would be an allocation on the phone's frame
+    /// path in exchange for nothing at these sizes.
+    fn remask_shape_users(
+        active: &HashMap<WidgetId, ActiveData>,
+        id: WidgetId,
+        own: MaskIdx,
+        freed: &[PrimitiveHandle],
+        rsc: &mut dyn UiRsc,
+    ) {
+        let mut stale: Vec<MaskIdx> = Vec::new();
+        for (i, mask) in rsc.ui().masks.iter().enumerate() {
+            let idx = Id::preset(i as u32);
+            if idx != own && freed.iter().any(|h| h.slot == mask.primitive) {
+                stale.push(idx);
+            }
+        }
+        if stale.is_empty() {
+            return;
+        }
+        let mut owners: Vec<WidgetId> = Vec::new();
+        for (widget, data) in active {
+            if *widget != id && stale.contains(&data.own_mask) {
+                owners.push(*widget);
+            }
+        }
+        for owner in owners {
+            rsc.widgets_mut().needs_redraw.insert(owner);
+        }
+    }
+
     fn remove_rec(&mut self, id: WidgetId, rsc: &mut dyn UiRsc) -> Option<ActiveData> {
         let inst = self.remove(id, true, rsc);
         if let Some(inst) = &inst {
@@ -970,6 +1037,92 @@ impl UiRenderState {
         parts.join(" -> ")
     }
 
+    /// One primitive's corners in window pixels -- the transliteration of
+    /// `shader.wgsl`'s `corners_of`, `floor` for `floor`. The rounding is
+    /// the whole reason this is not `region.to_px()`: the shader floors
+    /// each half separately before adding the move delta, and a hit test
+    /// that skipped it would disagree with the pixels by up to one along
+    /// each edge -- invisible in every test written against a whole-pixel
+    /// layout and wrong on the phone, whose 2.55 density makes nothing
+    /// land on a whole pixel.
+    pub fn primitive_corners(&self, slot: u32, rsc: &dyn UiRsc) -> PixelRegion {
+        let inst = self.primitives.instance(slot);
+        let delta = self.resolve_move_chain(inst.move_idx, rsc);
+        let size = self.output_size;
+        let corner = |c: UiVec2| (c.get_rel() * size).floor() + c.get_abs().floor() + delta;
+        PixelRegion {
+            top_left: corner(inst.region.top_left()),
+            bot_right: corner(inst.region.bot_right()),
+        }
+    }
+
+    /// Where a mask's clip actually is on screen: the box of the
+    /// primitive it references. Its *shape* within that box is
+    /// [`Self::mask_coverage`]'s -- this is the bounding box, which is
+    /// what a test asking "is the clip over the right part of the screen"
+    /// wants and all a square-cornered mask has ever had.
+    pub fn mask_region(&self, mask: MaskIdx, rsc: &dyn UiRsc) -> PixelRegion {
+        self.primitive_corners(rsc.ui().masks[mask.idx()].primitive, rsc)
+    }
+
+    /// How much of the pixel at `pos` (window pixels) survives `mask` and
+    /// every mask it nests inside: the referenced primitives' own
+    /// coverage, multiplied along the chain. The CPU half of
+    /// `shader.wgsl`'s `fs_main` mask loop -- same order, same bound, same
+    /// `rounded_rect_coverage` -- so a corner that cannot be tapped and a
+    /// corner that is not drawn are the same corner (LAYOUT.md's "Masks
+    /// with a shape", point 4).
+    ///
+    /// A mask whose shape is not a rect covers everything, exactly as the
+    /// shader's own `mask_coverage` does: `Painter::set_mask_to` rejects
+    /// those by name, so this is the unreachable half of the same
+    /// agreement rather than a second policy.
+    pub fn mask_coverage(&self, mask: MaskIdx, pos: Vec2, rsc: &dyn UiRsc) -> f32 {
+        let mut coverage = 1.0;
+        let mut at = mask;
+        for i in 0..PARENT_CHAIN_LIMIT {
+            if at == MaskIdx::NONE {
+                return coverage;
+            }
+            let m = rsc.ui().masks[at.idx()];
+            if let Some(rect) = self.primitives.primitive_data::<RectPrimitive>(m.primitive) {
+                let c = self.primitive_corners(m.primitive, rsc);
+                coverage *= rounded_rect_coverage(pos, c.top_left, c.bot_right, rect.radius);
+            }
+            at = m.parent;
+            debug_assert!(
+                i + 1 < PARENT_CHAIN_LIMIT || at == MaskIdx::NONE,
+                "mask chain exceeded PARENT_CHAIN_LIMIT ({PARENT_CHAIN_LIMIT}) from {mask:?} -- a \
+                 repeated slot means a `parent` link is cyclic, all-distinct slots mean the tree \
+                 nests deeper than shader.wgsl's own walk of the same bound",
+            );
+        }
+        coverage
+    }
+
+    /// Whether `pos` is inside `mask` at all -- more than half covered,
+    /// which is where the drawn edge is (`rounded_rect_coverage`'s doc).
+    /// What a hit test asks.
+    pub fn mask_admits(&self, mask: MaskIdx, pos: Vec2, rsc: &dyn UiRsc) -> bool {
+        self.mask_coverage(mask, pos, rsc) > 0.5
+    }
+
+    /// The first primitive `id`'s subtree wrote this frame, depth first
+    /// in draw order -- what a mask pointed at a widget clips to
+    /// (`Painter::set_mask_to_widget`). A widget that draws more than one
+    /// (a bordered rect is one primitive; a card with a stripe is two)
+    /// gives its first; a widget that wants another names it.
+    pub fn first_primitive(&self, id: WidgetId) -> Option<u32> {
+        let active = self.active.get(&id)?;
+        if let Some(h) = active.primitives.first() {
+            return Some(h.slot);
+        }
+        active
+            .children
+            .iter()
+            .find_map(|child| self.first_primitive(*child))
+    }
+
     pub fn window_region(&self, id: &impl IdLike, rsc: &dyn UiRsc) -> Option<PixelRegion> {
         let region = self.resolved_region(id, rsc)?;
         Some(region.to_px(self.output_size))

iris/macro/src/lib.rs

@@ -18,6 +18,12 @@ struct Input {
 }
 
 struct InputFn {
+    /// Everything written above the `fn` -- in practice a `///` doc
+    /// comment, which is why this exists: `masked_by` and its siblings
+    /// are public API and rustdoc is where their contract is read, so a
+    /// macro that silently rejected `///` sent the explanation into an
+    /// ordinary `//` comment nobody generating docs ever sees.
+    attrs: Vec<Attribute>,
     sig: Signature,
     body: Block,
 }
@@ -32,9 +38,10 @@ impl Parse for Input {
         input.parse::<Token![;]>()?;
         let mut fns = Vec::new();
         while !input.is_empty() {
+            let attrs = input.call(Attribute::parse_outer)?;
             let sig = input.parse()?;
             let body = input.parse()?;
-            fns.push(InputFn { sig, body })
+            fns.push(InputFn { attrs, sig, body })
         }
         if !input.is_empty() {
             input.error("function expected");
@@ -59,10 +66,15 @@ pub fn widget_trait(input: TokenStream) -> TokenStream {
         fns,
     } = parse_macro_input!(input as Input);
 
-    let sigs: Vec<_> = fns.iter().map(|f| f.sig.clone()).collect();
+    // The attributes go on the trait's own signature, which is the one
+    // rustdoc renders; the impl gets the bare `fn`.
+    let sigs: Vec<_> = fns
+        .iter()
+        .map(|InputFn { attrs, sig, .. }| quote! { #(#attrs)* #sig })
+        .collect();
     let impls: Vec<_> = fns
         .iter()
-        .map(|InputFn { sig, body }| quote! { #sig #body })
+        .map(|InputFn { sig, body, .. }| quote! { #sig #body })
         .collect();
 
     let Some(GenericParam::Type(state)) = generics.params.first() else {

iris/src/default/render.rs

@@ -82,19 +82,40 @@ impl UiRenderer {
     pub fn new(window: Arc<Window>) -> Self {
         let size = window.inner_size();
 
-        let instance = Instance::new(&InstanceDescriptor {
+        // `force-gles` on the desktop too, not just on Android: the
-            // `force-gles` on the desktop too, not just on Android: the
+        // GLES backend has behaviour of its own (a one-layer array
-            // GLES backend has behaviour of its own (a one-layer array
+        // texture is a `GL_TEXTURE_2D` -- see
-            // texture is a `GL_TEXTURE_2D` -- see
+        // `GpuTextures::create_array_texture`), and a machine with a
-            // `GpuTextures::create_array_texture`), and a machine with a
+        // real GPU is where that is cheap to reproduce and screenshot.
-            // real GPU is where that is cheap to reproduce and screenshot.
+        let mut backends = if cfg!(feature = "force-gles") {
-            backends: if cfg!(feature = "force-gles") {
+            Backends::GL
-                Backends::GL
+        } else {
-            } else {
+            Backends::PRIMARY
-                Backends::PRIMARY
+        };
-            },
+        let mut instance = Instance::new(&InstanceDescriptor {
+            backends,
             ..Default::default()
         });
+        // The same fallback the Android backend grew in 85869d0, and for
+        // the same reason: a machine can advertise a Vulkan ICD with no
+        // device behind it, and refusing to draw at all because the only
+        // usable adapter is a GLES one is iris's bug rather than the
+        // machine's. On this VM the virtio-gpu Venus device disappears
+        // whenever the host runs out of virgl contexts, so `run-headless.
+        // sh` -- layer 2 of the test rig -- aborted with `Could not get
+        // adapter!` while GL was sitting there working. Probed before the
+        // surface exists, matching Android, where an instance carrying
+        // both backends fails worse than one carrying the wrong one.
+        if backends != Backends::GL && instance.enumerate_adapters(backends).block_on().is_empty() {
+            log::warn!(
+                "iris renderer: no {backends:?} adapter on this machine, falling back to GLES"
+            );
+            backends = Backends::GL;
+            instance = Instance::new(&InstanceDescriptor {
+                backends,
+                ..Default::default()
+            });
+        }
 
         let surface = instance
             .create_surface(window.clone())
@@ -107,7 +128,9 @@ impl UiRenderer {
                 force_fallback_adapter: false,
             })
             .block_on()
-            .expect("Could not get adapter!");
+            .unwrap_or_else(|error| {
+                panic!("No usable GPU adapter for backends {backends:?}: {error}")
+            });
 
         // No features beyond what wgpu asks for by default, and no
         // binding-array limits: the atlas is one texture_2d_array and a

iris/src/layout_tests.rs

@@ -161,7 +161,10 @@ fn redrawing_a_masked_widget_does_not_nest_its_own_mask() {
         ui: UiData::default(),
     };
     let (_scroll, inner_root, _rects) = scrolled_rects(&mut rsc, 8);
-    let masked = rsc.ui.widgets.add_strong(Masked { inner: inner_root });
+    let masked = rsc.ui.widgets.add_strong(Masked {
+        shape: None,
+        inner: inner_root,
+    });
     let masked_id = masked.id();
     let root = masked.any();
     let mut render = UiRenderState::new();
@@ -184,7 +187,10 @@ fn a_mask_stays_put_while_its_scrolled_content_moves() {
         ui: UiData::default(),
     };
     let (scroll, inner_root, _rects) = scrolled_rects(&mut rsc, 500);
-    let masked = rsc.ui.widgets.add_strong(Masked { inner: inner_root });
+    let masked = rsc.ui.widgets.add_strong(Masked {
+        shape: None,
+        inner: inner_root,
+    });
     let masked_id = masked.id();
     let root = masked.any();
     let mut render = UiRenderState::new();
@@ -418,7 +424,10 @@ fn a_masked_widget_keeps_one_mask_slot_that_is_always_its_own_region() {
         ui: UiData::default(),
     };
     let (_scroll, inner_root, _rects) = scrolled_rects(&mut rsc, 8);
-    let masked = rsc.ui.widgets.add_strong(Masked { inner: inner_root });
+    let masked = rsc.ui.widgets.add_strong(Masked {
+        shape: None,
+        inner: inner_root,
+    });
     let masked_id = masked.id();
     // Placed at the bottom of a `Span::DOWN` behind a `rest(1)` sibling,
     // which is what moves the bar away from the provisional slot it is
@@ -453,7 +462,7 @@ fn a_masked_widget_keeps_one_mask_slot_that_is_always_its_own_region() {
     );
     let mask = *rsc.ui.masks.iter().next().unwrap();
     assert_eq!(
-        mask.region,
+        render.primitives.instance(mask.primitive).region,
         render.active.get(&masked_id).unwrap().region,
         "the mask a descendant clips against must be this widget's current box"
     );
@@ -655,3 +664,246 @@ fn a_widget_moved_by_its_parent_and_then_placed_inside_it_lands_at_the_placement
          after={after:?}"
     );
 }
+
+// ---------------------------------------------------------------------
+// LAYOUT.md's "Masks with a shape" -- its pass conditions, at layer 1.
+//
+// The shape a mask clips to is a *primitive already drawn*, never a copy
+// of one, so "the child's clipped corner" and "the container's own corner"
+// are the same arithmetic. These say so by evaluating both and demanding
+// exact equality: an approximate assertion would also pass a second copy
+// of the radius that merely happened to agree.
+// ---------------------------------------------------------------------
+
+const RADIUS: f32 = 20.0;
+
+/// A rounded container with `.masked_by` it, holding a `Rect::REST` child
+/// that fills it -- so the child's own corners are exactly the corners
+/// being clipped away. Returns the drawn state, the mask, the child, and
+/// the shape primitive the mask points at.
+fn rounded_container(rsc: &mut TestRsc) -> (UiRenderState, MaskIdx, WidgetId, u32) {
+    let child = rsc.ui.widgets.add_strong(Rect::new(UiColor::WHITE));
+    let child_id = child.id();
+    let shape = rsc
+        .ui
+        .widgets
+        .add_strong(Rect::new(UiColor::BLACK).radius(Len::abs(RADIUS)));
+    let shape_id = shape.id();
+    let root = rsc
+        .ui
+        .widgets
+        .add_strong(Masked {
+            shape: Some(shape.any()),
+            inner: child.any(),
+        })
+        .any();
+
+    let mut render = UiRenderState::new();
+    render.resize((200.0, 100.0));
+    render.update(&root, rsc);
+
+    let mask = render
+        .active
+        .get(&child_id)
+        .expect("the child is drawn")
+        .mask;
+    assert_ne!(
+        mask,
+        MaskIdx::NONE,
+        "the child was drawn with no clip at all"
+    );
+    let slot = render
+        .first_primitive(shape_id)
+        .expect("the shape widget drew a rect");
+    (render, mask, child_id, slot)
+}
+
+/// The pass condition: the child's coverage at a corner pixel *equals*
+/// the container's own coverage there. Exactly equal, because it is the
+/// same primitive evaluated once -- LAYOUT.md's point 1.
+#[test]
+fn a_masked_child_is_clipped_by_its_container_s_own_corner() {
+    let mut rsc = TestRsc {
+        ui: UiData::default(),
+    };
+    let (render, mask, _child, slot) = rounded_container(&mut rsc);
+    let corners = render.primitive_corners(slot, &rsc);
+    let radius = render
+        .primitives
+        .primitive_data::<RectPrimitive>(slot)
+        .expect("a mask's shape is a rect")
+        .radius;
+
+    // Across the whole corner arc, not one point on it: a single sample
+    // is satisfied by a mask that clips to the box and happens to agree
+    // where the two coincide. Swept from the arc's own centre -- the
+    // straight chord between the two ends of the arc lies *inside* the
+    // circle everywhere, so a walk along it never leaves the shape and
+    // the `outside` count below is what caught that.
+    let arc_center = corners.top_left + Vec2::new(radius, radius);
+    let (mut outside, mut inside) = (0, 0);
+    for i in 0..=20 {
+        let angle = std::f32::consts::FRAC_PI_2 * i as f32 / 20.0;
+        let dir = Vec2::new(-angle.cos(), -angle.sin());
+        for out in [-1.5f32, 0.0, 1.5] {
+            let pos = arc_center + dir * (radius + out);
+            let container = rounded_rect_coverage(pos, corners.top_left, corners.bot_right, radius);
+            assert_eq!(
+                render.mask_coverage(mask, pos, &rsc),
+                container,
+                "at {pos:?} the child's clip and the container's own edge disagree",
+            );
+            if container < 0.5 {
+                outside += 1;
+            } else {
+                inside += 1;
+            }
+        }
+    }
+    assert!(
+        outside > 0 && inside > 0,
+        "the sweep stayed on one side of the curve ({outside} out, {inside} in), so it proved \
+         nothing about the corner"
+    );
+}
+
+/// A hit test asks the same question the pixels do: the corner the
+/// container rounded away is not there to be pressed, and a point just
+/// inside the curve is. LAYOUT.md's point 4.
+#[test]
+fn a_mask_s_shape_decides_what_can_be_pressed() {
+    let mut rsc = TestRsc {
+        ui: UiData::default(),
+    };
+    let (render, mask, _child, slot) = rounded_container(&mut rsc);
+    let corners = render.primitive_corners(slot, &rsc);
+
+    // The very corner of the box, which the radius cut off.
+    let cut = corners.top_left + Vec2::new(1.0, 1.0);
+    assert!(
+        !render.mask_admits(mask, cut, &rsc),
+        "the corner the container rounded away is still pressable",
+    );
+    // The same distance in along the diagonal, past the curve.
+    let inside = corners.top_left + Vec2::new(RADIUS, RADIUS);
+    assert!(
+        render.mask_admits(mask, inside, &rsc),
+        "a point well inside the curve is not pressable",
+    );
+    // And the middle of an edge, which no radius touches -- the half the
+    // rounding had no reason to change.
+    let edge = Vec2::new(
+        (corners.top_left.x + corners.bot_right.x) / 2.0,
+        corners.top_left.y + 1.0,
+    );
+    assert!(
+        render.mask_admits(mask, edge, &rsc),
+        "a straight edge between two corners is not pressable",
+    );
+}
+
+/// Nested masks multiply, so a pixel inside two feathered corners is
+/// dimmed by both -- LAYOUT.md's point 2, and the "alpha should be
+/// decreased / multiplied" Iris asked for. Written as a product of the
+/// two the shader would compute separately, which is what "multiply"
+/// means and what an intersection test would get wrong.
+#[test]
+fn nested_masks_multiply_their_coverage() {
+    let mut rsc = TestRsc {
+        ui: UiData::default(),
+    };
+    let child = rsc.ui.widgets.add_strong(Rect::new(UiColor::WHITE));
+    let child_id = child.id();
+
+    let inner_shape = rsc
+        .ui
+        .widgets
+        .add_strong(Rect::new(UiColor::BLACK).radius(Len::abs(RADIUS)));
+    let inner_shape_id = inner_shape.id();
+    let inner = rsc.ui.widgets.add_strong(Masked {
+        shape: Some(inner_shape.any()),
+        inner: child.any(),
+    });
+    let outer_shape = rsc
+        .ui
+        .widgets
+        .add_strong(Rect::new(UiColor::BLACK).radius(Len::abs(RADIUS)));
+    let outer_shape_id = outer_shape.id();
+    let root = rsc
+        .ui
+        .widgets
+        .add_strong(Masked {
+            shape: Some(outer_shape.any()),
+            inner: inner.any(),
+        })
+        .any();
+
+    let mut render = UiRenderState::new();
+    render.resize((200.0, 100.0));
+    render.update(&root, &mut rsc);
+
+    let mask = render.active.get(&child_id).expect("drawn").mask;
+    let one = |render: &UiRenderState, rsc: &TestRsc, id, pos| {
+        let slot = render.first_primitive(id).expect("a shape rect");
+        let c = render.primitive_corners(slot, rsc);
+        let radius = render
+            .primitives
+            .primitive_data::<RectPrimitive>(slot)
+            .unwrap()
+            .radius;
+        rounded_rect_coverage(pos, c.top_left, c.bot_right, radius)
+    };
+
+    // A point on the corner arc, where both feathers are partial -- the
+    // only place a product and a minimum differ.
+    let slot = render.first_primitive(inner_shape_id).unwrap();
+    let corners = render.primitive_corners(slot, &rsc);
+    let pos = corners.top_left + Vec2::new(RADIUS * 0.3, RADIUS * 0.3);
+    let inner_cov = one(&render, &rsc, inner_shape_id, pos);
+    let outer_cov = one(&render, &rsc, outer_shape_id, pos);
+    assert!(
+        inner_cov > 0.0 && inner_cov < 1.0,
+        "the sample point is not inside a feather ({inner_cov}), so this proves nothing"
+    );
+    assert_eq!(
+        render.mask_coverage(mask, pos, &rsc),
+        inner_cov * outer_cov,
+        "two nested masks must multiply, not intersect",
+    );
+}
+
+/// A plain `.masked()` -- no shape given -- still clips to the widget's
+/// own box with square corners, which is what every list and scroll area
+/// relies on. The half the shape work had no reason to touch, and the one
+/// that would silently round every existing clip if `set_mask` ever wrote
+/// a radius of its own.
+#[test]
+fn a_plain_mask_still_clips_to_a_square_box() {
+    let mut rsc = TestRsc {
+        ui: UiData::default(),
+    };
+    let (_scroll, inner_root, _rects) = scrolled_rects(&mut rsc, 8);
+    let root = rsc
+        .ui
+        .widgets
+        .add_strong(Masked {
+            shape: None,
+            inner: inner_root,
+        })
+        .any();
+    let mut render = UiRenderState::new();
+    render.resize((200.0, 100.0));
+    render.update(&root, &mut rsc);
+
+    let mask = *rsc.ui.masks.iter().next().expect("one mask");
+    let corners = render.primitive_corners(mask.primitive, &rsc);
+    let mask_idx = MaskIdx::preset(0);
+    assert!(
+        render.mask_admits(mask_idx, corners.top_left + Vec2::new(0.5, 0.5), &rsc),
+        "a square clip must admit its own corner pixel",
+    );
+    assert!(
+        !render.mask_admits(mask_idx, corners.top_left - Vec2::new(2.0, 2.0), &rsc),
+        "a square clip must reject a point outside it",
+    );
+}

iris/src/sense.rs

@@ -292,7 +292,21 @@ impl SensorUi for UiRenderState {
             for (id, sensor) in active.get_mut(&layer).into_flat_iter() {
                 let shape = self.resolved_region(id, rsc).unwrap();
                 let region = shape.to_px(window_size);
-                let in_shape = cursor.exists && region.contains(cursor.pos);
+                // The mask this widget is drawn under, applied with the
+                // same coverage the fragment stage clips it with
+                // (LAYOUT.md's "Masks with a shape", point 4): a corner
+                // rounded away by a container is not there to be pressed,
+                // and a row scrolled out of a list's box is not either.
+                // The box test stays because it is what says the pointer
+                // is over *this* widget rather than merely inside its
+                // clip -- the two ask different questions and both have
+                // to hold.
+                let in_shape = cursor.exists
+                    && region.contains(cursor.pos)
+                    && self
+                        .active
+                        .get(id)
+                        .is_none_or(|a| self.mask_admits(a.mask, cursor.pos, rsc));
                 sensor.hover.update(in_shape);
                 if sensor.hover == ActivationState::Off {
                     continue;

iris/src/widget/list.rs

@@ -1315,6 +1315,7 @@ mod tests {
         let strong = rsc.ui.widgets.add_strong(list);
         let weak = strong.weak();
         let root = rsc.ui.widgets.add_strong(Masked {
+            shape: None,
             inner: strong.any(),
         });
         (weak, root.any())

iris/src/widget/mask.rs

@@ -1,12 +1,36 @@
 use crate::prelude::*;
 
+/// Clips `inner` -- and everything below it -- to a shape.
+///
+/// The shape is a **primitive**, never a rectangle or a radius stored
+/// here: with `shape`, the widget named there is drawn behind `inner`
+/// filling the same box and the clip is its first primitive, so a rounded
+/// container's corner and the corner its content is cut to are the same
+/// arithmetic and cannot fall out of step. Without one, this writes an
+/// undrawn rect at its own region, which is the plain "clip to my box"
+/// every list and scroll area wants. See docs/LAYOUT.md's "Masks with a
+/// shape".
 pub struct Masked {
+    /// The widget whose first primitive is the clip, drawn behind
+    /// `inner`, or `None` for this widget's own box.
+    pub shape: Option<StrongWidget>,
     pub inner: StrongWidget,
 }
 
 impl Widget for Masked {
     fn draw(&mut self, painter: &mut Painter) -> Size {
-        painter.set_mask(painter.region());
+        match &self.shape {
+            // Layered the way `Stack` layers a background under its
+            // content, and for the same reason: within one layer the draw
+            // order is undefined once anything has been freed.
+            Some(shape) => {
+                painter.child_layer();
+                painter.widget(shape);
+                painter.set_mask_to_widget(shape);
+                painter.next_layer();
+            }
+            None => painter.set_mask(painter.region()),
+        }
         painter.widget(&self.inner)
     }
 }

iris/src/widget/trait_fns.rs

@@ -87,11 +87,10 @@ widget_trait! {
         self.scrollable_on(Axis::Y)
     }
 
-    // `scrollable` along `axis`. A code fence pans across its own long
+    /// `scrollable` along `axis`. A code fence pans across its own long
-    // lines exactly the way a transcript pans down its rows, so the two
+    /// lines exactly the way a transcript pans down its rows, so the two
-    // are one function with the axis passed in rather than a second copy
+    /// are one function with the axis passed in rather than a second copy
-    // -- `DragArbiter::on` is the other half. (A `///` doc comment here
+    /// -- `DragArbiter::on` is the other half.
-    // is not accepted by `widget_trait!`, which parses its body itself.)
     fn scrollable_on(self, axis: Axis) -> impl WidgetIdFn<Rsc, Scroll> where Rsc: HasEvents {
         move |state| {
             Scroll::new(self.add_strong(state), axis)
@@ -119,6 +118,20 @@ widget_trait! {
 
     fn masked(self) -> impl WidgetFn<Rsc, Masked> {
         move |state| Masked {
+            shape: None,
+            inner: self.add_strong(state),
+        }
+    }
+
+    /// Clip to `shape` rather than to a plain box: `shape` is drawn
+    /// behind this widget, filling the same region, and what clips is the
+    /// primitive it drew -- so a rounded background and the corner its
+    /// content is cut to are one rect, with no radius passed twice.
+    /// Replaces `.masked().background(w)`, which drew the two but clipped
+    /// to the box.
+    fn masked_by<T>(self, shape: impl WidgetLike<Rsc, T>) -> impl WidgetFn<Rsc, Masked> {
+        move |state| Masked {
+            shape: Some(shape.add_strong(state)),
             inner: self.add_strong(state),
         }
     }

iris/tests/mask_sdf.rs

@@ -0,0 +1,247 @@
+//! The CPU rounded-rect SDF and the shader's own must agree.
+//!
+//! LAYOUT.md's "Masks with a shape" turns on it: the fragment stage clips
+//! a masked subtree with `shader.wgsl`'s `rounded_rect_coverage`, and the
+//! hit test (`UiRenderState::mask_admits`) clips the *same* subtree with
+//! `iris_core::rounded_rect_coverage`, so a corner that cannot be tapped
+//! and a corner that is not drawn are the same corner only while the two
+//! functions answer the same. Nothing else checks that: both sides are
+//! individually plausible and drift shows up as a control that is a pixel
+//! 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
+//! pass over a grid of points, and compares what came back with the Rust
+//! function at the same points. This is the only test in the workspace
+//! that needs a GPU; everything else about masks is layer 1 (docs/RUST.md's
+//! "Three test layers"). It fails rather than skips when there is no
+//! adapter, because a check that quietly did not run reads exactly like a
+//! check that passed.
+
+use iris_core::{SHAPE_SHADER, rounded_rect_coverage, util::Vec2};
+use pollster::FutureExt;
+use wgpu::util::DeviceExt;
+
+/// The rect the grid is sampled against, in window pixels. Deliberately
+/// off the whole-pixel grid: the shader floors a primitive's corners, but
+/// `rounded_rect_coverage` is handed pixels either side of that and has to
+/// agree at fractional positions too -- the phone's 2.55 density puts
+/// nothing on a whole pixel.
+const TOP_LEFT: Vec2 = Vec2::new(10.5, 20.25);
+const BOT_RIGHT: Vec2 = Vec2::new(170.75, 90.0);
+
+/// Radii spanning what the widgets actually ask for, plus the two edges of
+/// the function's own domain: a square corner, and one large enough that
+/// `min(edge, radius)` stops mattering.
+const RADII: [f32; 5] = [0.0, 0.75, 8.0, 20.0, 34.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
+/// [0, 1], so this is about six decimal digits -- four orders of magnitude
+/// tighter than the half-pixel feather the hit test reads, which is what
+/// the agreement is actually for.
+const TOLERANCE: f32 = 1e-5;
+
+#[test]
+fn mask_sdf_matches_the_shader() {
+    let points = grid();
+    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 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);
+        }
+    }
+    let (pos, radius, want, got) = worst_at;
+    assert!(
+        worst <= TOLERANCE,
+        "shader.wgsl's rounded_rect_coverage and iris_core's disagree by {worst} at {pos:?} \
+         (radius {radius}): the CPU says {want}, the GPU {got}. One of the two was edited \
+         without the other -- they are transliterations and have to stay so, or a masked \
+         corner stops being tappable where it is drawn.",
+    );
+
+    // 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), \
+         so agreeing proved nothing",
+    );
+}
+
+/// Points spanning the rect and a margin outside it, at half-pixel steps
+/// so the feather is sampled rather than stepped over, paired with each
+/// radius.
+fn grid() -> Vec<(Vec2, f32)> {
+    let mut points = Vec::new();
+    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.
+fn run_shader(points: &[(Vec2, f32)]) -> Vec<f32> {
+    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 (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 {
+        label: Some("mask sdf probe"),
+        layout: None,
+        module: &module,
+        entry_point: Some("probe"),
+        compilation_options: Default::default(),
+        cache: None,
+    });
+
+    // (x, y, radius, unused) -- one vec4 per point, so the buffer needs no
+    // stride arithmetic and no alignment rule of its own.
+    let input: Vec<[f32; 4]> = points
+        .iter()
+        .map(|(pos, radius)| [pos.x, pos.y, *radius, 0.0])
+        .collect();
+    let in_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+        label: Some("mask sdf points"),
+        contents: bytemuck::cast_slice(&input),
+        usage: wgpu::BufferUsages::STORAGE,
+    });
+    let out_size = (points.len() * std::mem::size_of::<f32>()) as u64;
+    let out_buf = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("mask sdf coverage"),
+        size: out_size,
+        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
+        mapped_at_creation: false,
+    });
+    let read_buf = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("mask sdf readback"),
+        size: out_size,
+        usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+        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());
+        pass.set_pipeline(&pipeline);
+        pass.set_bind_group(0, &bind_group, &[]);
+        pass.dispatch_workgroups(points.len().div_ceil(64) as u32, 1, 1);
+    }
+    enc.copy_buffer_to_buffer(&out_buf, 0, &read_buf, 0, out_size);
+    queue.submit([enc.finish()]);
+
+    let slice = read_buf.slice(..);
+    slice.map_async(wgpu::MapMode::Read, |r| r.expect("mapping the readback"));
+    device
+        .poll(wgpu::PollType::wait_indefinitely())
+        .expect("waiting for the probe");
+    let coverage = bytemuck::cast_slice::<u8, f32>(&slice.get_mapped_range()).to_vec();
+    read_buf.unmap();
+    coverage
+}
+
+/// 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
+/// shader is what this test is for, and a copy here would be edited along
+/// with it.
+fn probe_source() -> String {
+    format!(
+        "{}\n{}\n\
+         @group(0) @binding(0) var<storage, read> probe_in: array<vec4<f32>>;\n\
+         @group(0) @binding(1) var<storage, read_write> probe_out: array<f32>;\n\
+         @compute @workgroup_size(64)\n\
+         fn probe(@builtin(global_invocation_id) gid: vec3<u32>) {{\n\
+         let i = gid.x;\n\
+         if i >= arrayLength(&probe_out) {{ return; }}\n\
+         let p = probe_in[i];\n\
+         probe_out[i] = rounded_rect_coverage(p.xy, vec2({}, {}), vec2({}, {}), p.z);\n\
+         }}\n",
+        wgsl_fn("distance_from_rect"),
+        wgsl_fn("rounded_rect_coverage"),
+        TOP_LEFT.x,
+        TOP_LEFT.y,
+        BOT_RIGHT.x,
+        BOT_RIGHT.y,
+    )
+}
+
+/// One WGSL function's whole text, from its `fn` keyword to the `}` that
+/// closes its body, found by matching braces. Panics by name when the
+/// function is not there, which is what a rename looks like from here.
+fn wgsl_fn(name: &str) -> &'static str {
+    let start = SHAPE_SHADER
+        .find(&format!("fn {name}("))
+        .unwrap_or_else(|| panic!("shader.wgsl has no `fn {name}(` -- renamed, or gone"));
+    let body = SHAPE_SHADER[start..]
+        .find('{')
+        .expect("a wgsl fn signature is followed by its body");
+    let mut depth = 0usize;
+    for (i, c) in SHAPE_SHADER[start + body..].char_indices() {
+        match c {
+            '{' => depth += 1,
+            '}' => {
+                depth -= 1;
+                if depth == 0 {
+                    return &SHAPE_SHADER[start..start + body + i + 1];
+                }
+            }
+            _ => {}
+        }
+    }
+    panic!("`fn {name}`'s body in shader.wgsl is never closed");
+}

iris/transcript-fixture/tests/top_edge.rs

@@ -116,7 +116,7 @@ fn the_list_is_clipped_to_its_own_box() {
         active.mask != MaskIdx::NONE,
         "the transcript's list is drawn with nothing clipping it",
     );
-    let clip = h.rsc.ui.masks[active.mask.idx()].region.to_px(h.size());
+    let clip = h.render.mask_region(active.mask, &h.rsc);
     let list = list_box(&h, &screen);
     assert!(
         clip.top_left.y >= list.top_left.y - 0.5 && clip.bot_right.y <= list.bot_right.y + 0.5,

iris/transcript-ui/src/row.rs

@@ -266,11 +266,14 @@ where
     // delta path knowing anything about frames.
     let framed = match frame {
         BlockFrame::Plain => field.width(rest(1)).add_strong(rsc).any(),
+        // Masked *by the panel*, not inside it: the fence's own rounded
+        // rect is the clip, so content scrolled sideways is cut on the
+        // curve instead of leaving square pixels in the corners
+        // (Iris, 2026-09-07).
         BlockFrame::Verbatim { fill } => field
             .scrollable_on(Axis::X)
-            .masked()
             .pad(dp(FRAME_PAD_DP))
-            .background(rect(fill).radius(dp(FRAME_RADIUS_DP)))
+            .masked_by(rect(fill).radius(dp(FRAME_RADIUS_DP)))
             .width(rest(1))
             .add_strong(rsc)
             .any(),