diff --git a/docs/IRIS.md b/docs/IRIS.md index f634dee..6e470f7 100644 --- a/docs/IRIS.md +++ b/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 diff --git a/docs/IRIS_TODO.md b/docs/IRIS_TODO.md index e4dc484..46a06e0 100644 --- a/docs/IRIS_TODO.md +++ b/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 - by something *and also* applies mask B — a mask can reference a parent - mask, the way the move chain references a parent offset. Today masks - are independent regions. Design it beside the move chain (same shape: - a parent index and a bounded walk in the shader); do it when a real - widget needs it, not before. +- [x] **Masks defined relative to each other. (Done: chaining + 2026-09-07 in d507ae4, the multiply 2026-09-08.)** Built exactly + beside the move chain, as this asked: `Mask::parent` is a slot index + and the fragment stage walks it under the same bound the move chain + uses. Each step multiplies the referenced primitive's coverage into + 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 diff --git a/docs/LAYOUT.md b/docs/LAYOUT.md index 83db8d2..257f7ec 100644 --- a/docs/LAYOUT.md +++ b/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. diff --git a/docs/RUST.md b/docs/RUST.md index c9d28ca..e8e8afb 100644 --- a/docs/RUST.md +++ b/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 - insets shown in the diagnostics overlay. The worktree note here was - stale by 2026-09-07 night: no worktree exists and the keyboard half - is ticked in IRIS_TODO's night entry. What remains is the impulse - estimator item below. +- [x] Fling parity with Compose, and the phone's keyboard push-up, with + insets shown in the diagnostics overlay. Every part is ticked in + IRIS_TODO's 2026-09-07 entries: the keyboard half in the night entry, + the estimator in the "later" one (Lsq2, not Impulse -- see the box + 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 - estimator with min/max fling velocity (docs/IRIS_TODO.md, 2026-09-07 - later). After the culling fix lands (same file). +- [x] Scroll clamped at both ends (e922b73, `List::clamp_to_content`) + and Compose's velocity estimator (docs/IRIS_TODO.md, 2026-09-07 + 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 - with no slop (docs/IRIS_TODO.md, night). Opus, next slot; uses the - layer-1 harness. +- [x] **Catch-a-fling (done 2026-09-07, b87f5a5; verified and ticked + 2026-09-08).** A down on a moving list ends the fling on that sample + 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 - 2026-09-07)". A mask references a primitive already drawn - (rect SDF, texture or glyph alpha), chained and multiplied; `.masked()` - points at the widget's own primitives; hit-testing applies the shape. - **Chaining landed early**, 2026-09-07 (d507ae4): a mask carries the - 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 - rectangles only -- the shape half, and the hit-testing half, are what - is left of this item. +- [x] **Masks with a shape (done 2026-09-08).** docs/LAYOUT.md's "Masks + with a shape" carries the design and, at its end, the four places the + code is deliberately narrower than it. **Chaining landed early**, + 2026-09-07 (d507ae4). What landed now is the shape half and the + hit-testing half: + + `Mask` is `{ primitive, parent }` -- the slot of a primitive already + written, and the mask this one nests inside. The fragment stage + 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 diff --git a/iris/Cargo.lock b/iris/Cargo.lock index 00d8d8f..aab4bef 100644 --- a/iris/Cargo.lock +++ b/iris/Cargo.lock @@ -1737,6 +1737,7 @@ dependencies = [ "accesskit_winit", "android-view", "arboard", + "bytemuck", "image", "iris-core", "iris-macro", diff --git a/iris/Cargo.toml b/iris/Cargo.toml index f747e52..00715ec 100644 --- a/iris/Cargo.toml +++ b/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]` diff --git a/iris/core/src/render/data.rs b/iris/core/src/render/data.rs index eed375b..aa130fe 100644 --- a/iris/core/src/render/data.rs +++ b/iris/core/src/render/data.rs @@ -49,22 +49,47 @@ impl MaskIdx { pub type MoveIdx = Id; +/// 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 mask-owning widget's own move slot -- resolved in the fragment - /// shader against the same chain the vertex shader walks for a - /// primitive's own corners, so a mask and the content clipped by it - /// can move independently. See LAYOUT.md section 2b. - pub move_idx: MoveIdx, + /// The slot in `UiRenderState::primitives` of the primitive whose + /// coverage this mask is. Today always a `RectPrimitive`: a glyph or + /// a standalone image would need, respectively, a CPU-side alpha + /// plane for the hit test to agree with the shader, and a bind-group + /// switch the fragment stage cannot make -- `Painter::set_mask` + /// 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 - /// to be inside every mask on it. Chained rather than intersected on - /// the CPU because each mask moves with its own widget -- a code fence - /// inside a transcript row carries the row's scroll, the list's 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. + /// clipping nests: the fragment stage walks the chain and multiplies + /// every coverage on it, which is what makes a pixel inside two + /// feathered corners dimmed by both. Chained rather than intersected + /// on the CPU because each mask moves with its own widget -- a code + /// fence inside a transcript row carries the row's scroll, the list's + /// 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 diff --git a/iris/core/src/render/mod.rs b/iris/core/src/render/mod.rs index 5452c97..92b75ed 100644 --- a/iris/core/src/render/mod.rs +++ b/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 diff --git a/iris/core/src/render/primitive.rs b/iris/core/src/render/primitive.rs index 3c81c01..8230517 100644 --- a/iris/core/src/render/primitive.rs +++ b/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; + /// 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; } macro_rules! primitives { @@ -86,6 +90,9 @@ macro_rules! primitives { fn vec(data: &mut PrimitiveData) -> &mut PrimitiveVec { &mut data.$name } + fn vec_ref(data: &PrimitiveData) -> &PrimitiveVec { + &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(&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 -/// layer position (`pos == NOT_DRAWN`) is a primitive that exists to be -/// *referenced* -- a mask's shape -- and is never rasterized. +/// where in a layer's draw order it currently sits ([`NOT_DRAWN`] if it is +/// only referenced). #[derive(Debug)] pub struct PrimitiveHandle { pub layer: usize, diff --git a/iris/core/src/render/sdf.rs b/iris/core/src/render/sdf.rs new file mode 100644 index 0000000..64a11ce --- /dev/null +++ b/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) +} diff --git a/iris/core/src/render/shader.wgsl b/iris/core/src/render/shader.wgsl index 0b6a737..fe15b96 100644 --- a/iris/core/src/render/shader.wgsl +++ b/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, - 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. + primitive: u32, parent: u32, } @@ -57,11 +55,6 @@ struct UiScalar { abs: f32, } -struct UiVec2 { - rel: vec2, - abs: vec2, -} - // 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> @@ -157,6 +150,27 @@ struct Region { bot_right: vec2, } +/// 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, + bot_right: vec2, +} + +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 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); - 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 c = corners_of(inst); + let top_left = c.top_left; + let bot_right = c.bot_right; let size = bot_right - top_left; let uv = vec2( @@ -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 - // own move slot (`Mask::parent` in data.rs). + // its own mask inside another is clipped by both, and the coverages + // 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); - 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; - } + color.a *= mask_coverage(pos, mask); 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, 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 { return textureSample(image_texture, samp, region.uv); } @@ -251,19 +270,35 @@ fn draw_glyph(region: Region, g: GlyphInfo) -> vec4 { 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, + top_left: vec2, + bot_right: vec2, + 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 { var color = unpack4x8unorm(rect.color); let edge = 0.5; - let size = region.bot_right - region.top_left; - 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); + color.a *= rounded_rect_coverage(region.pos, region.top_left, region.bot_right, rect.radius); 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); } diff --git a/iris/core/src/ui/painter.rs b/iris/core/src/ui/painter.rs index 7912780..b2091f4 100644 --- a/iris/core/src/ui/painter.rs +++ b/iris/core/src/ui/painter.rs @@ -1,7 +1,10 @@ use crate::{ - RenderedText, Size, StrongWidget, TextAttrs, TextBuffer, TextData, TextureHandle, UiRegion, - UiRenderState, UiRsc, UiScalar, UiVec2, WidgetId, - render::{GlyphPrimitive, Mask, MaskIdx, MoveIdx, Primitive, PrimitiveHandle, PrimitiveInst}, + Color, RenderedText, Size, StrongWidget, TextAttrs, TextBuffer, TextData, TextureHandle, + UiRegion, UiRenderState, UiRsc, UiScalar, UiVec2, WidgetId, + 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(&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( + &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 - /// inside both, which is what lets a transcript row's code fence clip + /// one (`Mask::parent`) and the fragment stage multiplies both + /// 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(&mut self, shape: &StrongWidget) { + 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, - move_idx: self.move_slot, + primitive: shape, parent, }; let old_parent = if self.own_mask == MaskIdx::NONE { diff --git a/iris/core/src/ui/render_state.rs b/iris/core/src/ui/render_state.rs index ce3cae8..d83976a 100644 --- a/iris/core/src/ui/render_state.rs +++ b/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( &mut self, layer: usize, + drawn: Drawn, inst: PrimitiveInst

, ) -> 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, + id: WidgetId, + own: MaskIdx, + freed: &[PrimitiveHandle], + rsc: &mut dyn UiRsc, + ) { + let mut stale: Vec = 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 = 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 { 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::(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 { + 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 { let region = self.resolved_region(id, rsc)?; Some(region.to_px(self.output_size)) diff --git a/iris/macro/src/lib.rs b/iris/macro/src/lib.rs index 98b13cf..6263369 100644 --- a/iris/macro/src/lib.rs +++ b/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, sig: Signature, body: Block, } @@ -32,9 +38,10 @@ impl Parse for Input { input.parse::()?; 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 { diff --git a/iris/src/default/render.rs b/iris/src/default/render.rs index 9cac43d..ac4132f 100644 --- a/iris/src/default/render.rs +++ b/iris/src/default/render.rs @@ -82,19 +82,40 @@ impl UiRenderer { pub fn new(window: Arc) -> Self { let size = window.inner_size(); - let instance = Instance::new(&InstanceDescriptor { - // `force-gles` on the desktop too, not just on Android: the - // GLES backend has behaviour of its own (a one-layer array - // texture is a `GL_TEXTURE_2D` -- see - // `GpuTextures::create_array_texture`), and a machine with a - // real GPU is where that is cheap to reproduce and screenshot. - backends: if cfg!(feature = "force-gles") { - Backends::GL - } else { - Backends::PRIMARY - }, + // `force-gles` on the desktop too, not just on Android: the + // GLES backend has behaviour of its own (a one-layer array + // texture is a `GL_TEXTURE_2D` -- see + // `GpuTextures::create_array_texture`), and a machine with a + // real GPU is where that is cheap to reproduce and screenshot. + let mut backends = if cfg!(feature = "force-gles") { + Backends::GL + } else { + 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 diff --git a/iris/src/layout_tests.rs b/iris/src/layout_tests.rs index baa67e3..9e54d3c 100644 --- a/iris/src/layout_tests.rs +++ b/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::(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::(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", + ); +} diff --git a/iris/src/sense.rs b/iris/src/sense.rs index 3598ada..be2ee57 100644 --- a/iris/src/sense.rs +++ b/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; diff --git a/iris/src/widget/list.rs b/iris/src/widget/list.rs index 02e07bc..5a67e57 100644 --- a/iris/src/widget/list.rs +++ b/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()) diff --git a/iris/src/widget/mask.rs b/iris/src/widget/mask.rs index 298e052..1373599 100644 --- a/iris/src/widget/mask.rs +++ b/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, 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) } } diff --git a/iris/src/widget/trait_fns.rs b/iris/src/widget/trait_fns.rs index 0d3a98a..89eea22 100644 --- a/iris/src/widget/trait_fns.rs +++ b/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 - // 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 - // -- `DragArbiter::on` is the other half. (A `///` doc comment here - // is not accepted by `widget_trait!`, which parses its body itself.) + /// `scrollable` along `axis`. A code fence pans across its own long + /// 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 + /// -- `DragArbiter::on` is the other half. fn scrollable_on(self, axis: Axis) -> impl WidgetIdFn where Rsc: HasEvents { move |state| { Scroll::new(self.add_strong(state), axis) @@ -119,6 +118,20 @@ widget_trait! { fn masked(self) -> impl WidgetFn { 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(self, shape: impl WidgetLike) -> impl WidgetFn { + move |state| Masked { + shape: Some(shape.add_strong(state)), inner: self.add_strong(state), } } diff --git a/iris/tests/mask_sdf.rs b/iris/tests/mask_sdf.rs new file mode 100644 index 0000000..6201ae1 --- /dev/null +++ b/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 { + 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::()) 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::(&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 probe_in: array>;\n\ + @group(0) @binding(1) var probe_out: array;\n\ + @compute @workgroup_size(64)\n\ + fn probe(@builtin(global_invocation_id) gid: vec3) {{\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"); +} diff --git a/iris/transcript-fixture/tests/top_edge.rs b/iris/transcript-fixture/tests/top_edge.rs index 3589cb3..d1d34b7 100644 --- a/iris/transcript-fixture/tests/top_edge.rs +++ b/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, diff --git a/iris/transcript-ui/src/row.rs b/iris/transcript-ui/src/row.rs index 764ff07..7823f8d 100644 --- a/iris/transcript-ui/src/row.rs +++ b/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(),