iris is the framework alone; the app is one crate in app-rust/

Iris: "the organization of the rust rewrite is a mess right now... there
shouldn't be anything related to the app inside of iris. Iris is supposed
to be the UI framework alone." And, on the crate count: "I'm confused why
the app only code needs more than one crate though."

Nine cargo workspaces become three, and the port's project code -- which
sat in five places, four of them inside the framework -- becomes one crate,
`ai-app`, in `app-rust/`:

  client-core                -> app-rust/src/client
  iris/transcript-ui         -> app-rust/src/ui
  iris/transcript-fixture    -> app-rust/src/ui/fixture.rs + tests/ + touch/
  iris/desktop-app           -> app-rust/src/desktop + src/bin_desktop.rs
  iris/android-app           -> app-rust/src/android + android-project/
  android-shell              -> app-rust/src/shell

iris/ keeps core, macro, the iris crate, tabs-ui and rig-input, and now
mentions no session, transcript, setup or server anywhere.

Only two of the old splits had a reason that survived reading. event-model
stays a crate at the repo root because server/ depends on it too, so a
crate is what makes the backend and the app agree by construction. The two
Android .so names looked like a hard constraint -- a package produces one
library artifact -- until P2 turned out to already plan merging those two
Android apps into one; both faces now come out of libai_app.so, picked
apart by features so `--no-default-features --features shell` keeps wgpu,
parley and iris out of the Compose app's APK. docs/RUST.md's "One app
crate" has the rest, including what each remaining feature is for.

DECISIONS.md and SUBAGENTS.md move into docs/ with everything else.

Verified: ./run-tests.sh and `cd iris && cargo test` green, clippy and fmt
clean in all five workspaces, `cargo ndk -t x86_64` links libai_app.so,
build-apk.sh produces an APK that installs and launches on this checkout's
emulator (Gl ... virgl, as expected), and the phone-sized headless
screenshot renders the transcript unchanged.

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

core/src/attr.rs

@@ -1,4 +1,4 @@
-use crate::{UiRsc, WidgetIdFn, WidgetLike, WeakWidget};
+use crate::{UiRsc, WeakWidget, WidgetIdFn, WidgetLike};
 
 pub trait WidgetAttr<Rsc, W: ?Sized> {
     type Input;

core/src/event/manager.rs

@@ -79,6 +79,8 @@ type EventData<Rsc, E> = (E, Rc<dyn for<'a> EventFn<Rsc, <E as Event>::Data<'a>>
 pub struct TypeEventManager<Rsc: HasEvents, E: Event> {
     // TODO: reduce visiblity!!
     pub active: HashMap<LayerId, HashMap<WidgetId, E::State>>,
+    /// This event's own input-wide state -- see [`Event::Global`].
+    pub global: E::Global,
     map: HashMap<WidgetId, Vec<EventData<Rsc, E>>>,
 }
 
@@ -107,6 +109,7 @@ impl<Rsc: HasEvents, E: Event> Default for TypeEventManager<Rsc, E> {
     fn default() -> Self {
         Self {
             active: Default::default(),
+            global: Default::default(),
             map: Default::default(),
         }
     }
@@ -135,6 +138,18 @@ impl<Rsc: HasEvents + 'static, E: Event> TypeEventManager<Rsc, E> {
         ));
     }
 
+    /// The event lists this widget was registered with (`register`'s
+    /// `event` argument, one per call), without running anything. Lets a
+    /// caller ask "would this widget's registrations match the current
+    /// state" separately from actually dispatching to it -- used by
+    /// `sense.rs` to decide whether a widget genuinely consumes a scroll
+    /// or press this frame (so a lower layer can still receive it if not)
+    /// without that decision being conflated with "the cursor happens to
+    /// be over it," which is all `run_fn` running something tells you.
+    pub fn registered(&self, id: WidgetId) -> impl Iterator<Item = &E> {
+        self.map.get(&id).into_iter().flatten().map(|(e, _)| e)
+    }
+
     pub fn run_fn<'a>(
         &mut self,
         id: impl IdLike,

core/src/event/mod.rs

@@ -9,6 +9,20 @@ pub use rsc::*;
 pub trait Event: Sized + 'static + Clone {
     type Data<'a>: Clone = ();
     type State: Default = ();
+    /// State this event owns that belongs to no single widget -- what the
+    /// thing dispatching the event knows about the *input*, rather than
+    /// about a listener. `()` for almost every event; the cursor's is
+    /// `iris::sense::PointerInput` (which widget holds pointer capture,
+    /// and who is tracking the press in flight).
+    ///
+    /// It lives here so that such state has one owner, reached by `&mut`
+    /// through the event manager, instead of being parked on whatever
+    /// structure a handler happens to be able to reach and guarded with a
+    /// lock. Iris asked for that on 2026-09-08, of the pointer capture
+    /// that used to sit in a `Mutex` on `UiRenderState`: "everything
+    /// global should be stored in the general input handler, not in
+    /// specific senses with locking stuff."
+    type Global: Default = ();
     #[allow(unused_variables)]
     fn should_run<'a>(&self, data: &Self::Data<'a>) -> Option<Self::Data<'a>> {
         Some(data.clone())

core/src/icon.rs

@@ -0,0 +1,39 @@
+//! The icons iris draws, as codepoints in the Nerd Fonts subset it ships.
+//!
+//! **Why a bundled font rather than ordinary Unicode**: the disclosure
+//! mark used to be U+25B8/25BE/25B4 out of whatever face the platform
+//! resolved, and once iris stopped bundling fonts (DECISIONS.md,
+//! 2026-09-07) Iris's phone drew an empty box for them and this VM drew a
+//! dot. UI_RULES' answer is not to avoid glyphs but to ship them, which is
+//! also what the Compose app has always done for its icons
+//! (`app/build-icon-font.sh`, `NerdIcons.kt`) -- the same Material Design
+//! family, so an icon means the same thing in both apps.
+//!
+//! **Why not vector assets or drawn shapes**: an icon beside a line of
+//! text wants that line's size, colour and baseline, and text gets all
+//! three for free. This replaced `iris::widget::mark`, which drew the
+//! triangle into a texture: correct, but one shape, and every further icon
+//! would have been another bespoke rasteriser.
+//!
+//! Each constant here has to have a matching codepoint in
+//! `iris/core/build-icon-font.sh`'s `GLYPHS`; a codepoint here that the
+//! script did not subset is a glyph that silently isn't there. The subset
+//! is the font's **Mono** face, where every glyph is one em wide and one
+//! em tall, so two icons at one font size are one size without either
+//! being given one -- and why an icon looks smaller than text at the same
+//! size, since the glyph is drawn inside that em rather than filling it.
+//!
+//! Draw one with [`crate::Family::Icons`]:
+//!
+//! ```ignore
+//! text(icon::OPEN, 12.0, MUTED).family(Family::Icons)
+//! ```
+
+/// `md-menu_down` -- a filled triangle pointing down: this card is open.
+pub const OPEN: &str = "\u{F035D}";
+
+/// `md-menu_right` -- pointing right: this card opens.
+pub const CLOSED: &str = "\u{F035F}";
+
+/// `md-menu_up` -- pointing up: fold this group of cards away again.
+pub const COLLAPSE: &str = "\u{F0360}";

core/src/lib.rs

@@ -2,12 +2,9 @@
 #![feature(const_ops)]
 #![feature(const_trait_impl)]
 #![feature(const_convert)]
-#![feature(map_try_insert)]
 #![feature(unboxed_closures)]
 #![feature(fn_traits)]
-#![feature(const_cmp)]
 #![feature(const_destruct)]
-#![feature(portable_simd)]
 #![feature(associated_type_defaults)]
 #![feature(unsize)]
 #![feature(coerce_unsized)]
@@ -22,6 +19,7 @@ mod render;
 mod ui;
 mod widget;
 
+pub mod icon;
 pub mod util;
 
 pub use attr::*;

core/src/num.rs

@@ -5,19 +5,19 @@ pub const trait UiNum {
     fn to_f32(self) -> f32;
 }
 
-impl const UiNum for f32 {
+const impl UiNum for f32 {
     fn to_f32(self) -> f32 {
         self
     }
 }
 
-impl const UiNum for u32 {
+const impl UiNum for u32 {
     fn to_f32(self) -> f32 {
         self as f32
     }
 }
 
-impl const UiNum for i32 {
+const impl UiNum for i32 {
     fn to_f32(self) -> f32 {
         self as f32
     }
@@ -27,7 +27,7 @@ pub const fn vec2(x: impl const UiNum, y: impl const UiNum) -> Vec2 {
     Vec2::new(x.to_f32(), y.to_f32())
 }
 
-impl<T: const UiNum + Copy> const From<T> for Vec2 {
+const impl<T: const UiNum + Copy> From<T> for Vec2 {
     fn from(v: T) -> Self {
         Self {
             x: v.to_f32(),
@@ -36,7 +36,7 @@ impl<T: const UiNum + Copy> const From<T> for Vec2 {
     }
 }
 
-impl<T: const UiNum, U: const UiNum> const From<(T, U)> for Vec2
+const impl<T: const UiNum, U: const UiNum> From<(T, U)> for Vec2
 where
     (T, U): const Destruct,
 {

core/src/orientation/align.rs

@@ -187,7 +187,7 @@ impl From<CardinalAlign> for Align {
     }
 }
 
-impl const From<RegionAlign> for UiVec2 {
+const impl From<RegionAlign> for UiVec2 {
     fn from(align: RegionAlign) -> Self {
         Self::rel(align.rel())
     }

core/src/orientation/axis.rs

@@ -1,6 +1,6 @@
 use super::*;
 
-#[derive(Copy, Clone, Eq, PartialEq)]
+#[derive(Copy, Clone, Eq, PartialEq, Debug)]
 pub enum Axis {
     X,
     Y,
@@ -74,14 +74,14 @@ pub const trait AxisT {
 }
 
 pub struct XAxis;
-impl const AxisT for XAxis {
+const impl AxisT for XAxis {
     fn get() -> Axis {
         Axis::X
     }
 }
 
 pub struct YAxis;
-impl const AxisT for YAxis {
+const impl AxisT for YAxis {
     fn get() -> Axis {
         Axis::Y
     }

core/src/orientation/len.rs

@@ -9,7 +9,31 @@ pub struct Size {
 
 #[derive(Debug, Clone, Copy, PartialEq)]
 pub struct Len {
+    /// Physical pixels -- a raw device pixel, unaffected by the display's
+    /// density. Rare to want directly (a hairline border is the usual
+    /// case); most sizes should be `dp` instead. See `dp`'s own doc for why
+    /// the two are kept separate rather than one field a caller has to
+    /// remember to pre-multiply.
     pub abs: f32,
+    /// Density-independent pixels -- Android's `dp` / CSS's reference pixel
+    /// (1 unit = 1/160in), resolved against the display's density at
+    /// layout time (`apply_rest`'s `density` parameter) rather than at the
+    /// point a widget is built, since density is a property of the device
+    /// this ends up running on, not of the widget tree. This is the unit
+    /// IRIS_TODO.md's "a density-independent length unit" item asked for,
+    /// 2026-09-06: before it existed, every size in the tree was `abs`
+    /// (physical pixels), and the only way to make a 16px design draw at
+    /// the right *size* on a denser display was a single global multiply
+    /// applied to the whole rendered scene after layout -- which is also
+    /// what made text blurry (RUST.md's P0 box, "blurry ... glyphs drawn
+    /// at logical size and stretched by the scale"): a glyph rasterised at
+    /// 16 physical px and then stretched 3x by that global multiply is a
+    /// 48px area sampled from a 16px bitmap. Resolving `dp` per-length at
+    /// layout time instead means the font size handed to the text shaper
+    /// is already the physical size (`16.0.dp() * 3.0`), so the glyph
+    /// atlas rasterises at the display's real resolution and nothing
+    /// downstream needs to stretch anything.
+    pub dp: f32,
     pub rel: f32,
     pub rest: f32,
 }
@@ -67,10 +91,10 @@ impl Size {
         }
     }
 
-    pub fn to_uivec2(self) -> UiVec2 {
+    pub fn to_uivec2(self, density: f32) -> UiVec2 {
         UiVec2 {
-            x: self.x.apply_rest(),
-            y: self.y.apply_rest(),
+            x: self.x.apply_rest(density),
+            y: self.y.apply_rest(density),
         }
     }
 
@@ -98,26 +122,66 @@ impl Size {
 impl Len {
     pub const ZERO: Self = Self {
         abs: 0.0,
+        dp: 0.0,
         rel: 0.0,
         rest: 0.0,
     };
 
     pub const REST: Self = Self {
         abs: 0.0,
+        dp: 0.0,
         rel: 0.0,
         rest: 1.0,
     };
 
-    pub fn apply_rest(&self) -> UiScalar {
+    /// Resolves to a `UiScalar`, folding `dp` into `abs` pixels against
+    /// `density` (physical pixels per dp -- 1.0 on a desktop or an
+    /// unscaled display, `content_scale` on Android; see `dp`'s field
+    /// doc). Every other component of `Len` is already resolution-
+    /// independent (`rel` is a fraction of the parent; `rest` becomes a
+    /// fraction too, below), so `density` only ever touches this one term.
+    pub fn apply_rest(&self, density: f32) -> UiScalar {
         UiScalar {
             rel: self.rel + if self.rest > 0.0 { 1.0 } else { 0.0 },
-            abs: self.abs,
+            abs: self.abs + self.dp * density,
+        }
+    }
+
+    /// The same fold as [`Self::apply_rest`] but staying a `Len`, so
+    /// `rest` survives: `dp` becomes physical pixels and every other
+    /// component is left alone.
+    ///
+    /// **A `Len` a widget *reports* must have been through this.** `dp` is
+    /// an input unit -- a number the widget author wrote -- and the
+    /// containers that consume a reported length read `abs`/`rel`/`rest`
+    /// directly (`Span::draw`'s placement arithmetic, `Pad`'s addition),
+    /// so a reported `dp` is silently worth zero. That is what made the
+    /// composer's bar collapse to nothing the moment its content grew past
+    /// `MaxSize`'s cap: the cap was `dp(168)` and was returned unresolved,
+    /// so the bar was given a slot of 0 and the field inside it was panned
+    /// out of a container measured at -63px. `UiRenderState::draw_inner`
+    /// debug-asserts the invariant after every `Widget::draw`.
+    pub fn fold_dp(&self, density: f32) -> Self {
+        Self {
+            abs: self.abs + self.dp * density,
+            dp: 0.0,
+            rel: self.rel,
+            rest: self.rest,
         }
     }
 
     pub fn abs(abs: impl UiNum) -> Self {
         Self {
             abs: abs.to_f32(),
+            dp: 0.0,
+            rel: 0.0,
+            rest: 0.0,
+        }
+    }
+    pub fn dp(dp: impl UiNum) -> Self {
+        Self {
+            abs: 0.0,
+            dp: dp.to_f32(),
             rel: 0.0,
             rest: 0.0,
         }
@@ -125,6 +189,7 @@ impl Len {
     pub fn rel(rel: impl UiNum) -> Self {
         Self {
             abs: 0.0,
+            dp: 0.0,
             rel: rel.to_f32(),
             rest: 0.0,
         }
@@ -132,6 +197,7 @@ impl Len {
     pub fn rest(ratio: impl UiNum) -> Self {
         Self {
             abs: 0.0,
+            dp: 0.0,
             rel: 0.0,
             rest: ratio.to_f32(),
         }
@@ -144,6 +210,15 @@ pub mod len_fns {
     pub fn abs(abs: impl UiNum) -> Len {
         Len {
             abs: abs.to_f32(),
+            dp: 0.0,
+            rel: 0.0,
+            rest: 0.0,
+        }
+    }
+    pub fn dp(dp: impl UiNum) -> Len {
+        Len {
+            abs: 0.0,
+            dp: dp.to_f32(),
             rel: 0.0,
             rest: 0.0,
         }
@@ -151,6 +226,7 @@ pub mod len_fns {
     pub fn rel(rel: impl UiNum) -> Len {
         Len {
             abs: 0.0,
+            dp: 0.0,
             rel: rel.to_f32(),
             rest: 0.0,
         }
@@ -158,14 +234,15 @@ pub mod len_fns {
     pub fn rest(ratio: impl UiNum) -> Len {
         Len {
             abs: 0.0,
+            dp: 0.0,
             rel: 0.0,
             rest: ratio.to_f32(),
         }
     }
 }
 
-impl_op!(Len Add add; abs rel rest);
-impl_op!(Len Sub sub; abs rel rest);
+impl_op!(Len Add add; abs dp rel rest);
+impl_op!(Len Sub sub; abs dp rel rest);
 
 impl_op!(Size Add add; x y);
 impl_op!(Size Sub sub; x y);
@@ -187,6 +264,9 @@ impl std::fmt::Display for Len {
         if self.abs != 0.0 {
             write!(f, "{} abs;", self.abs)?;
         }
+        if self.dp != 0.0 {
+            write!(f, "{} dp;", self.dp)?;
+        }
         if self.rel != 0.0 {
             write!(f, "{} rel;", self.rel)?;
         }

core/src/orientation/pos.rs

@@ -124,13 +124,13 @@ impl Display for UiVec2 {
 impl_op!(UiVec2 Add add; x y);
 impl_op!(UiVec2 Sub sub; x y);
 
-impl const From<Vec2> for UiVec2 {
+const impl From<Vec2> for UiVec2 {
     fn from(abs: Vec2) -> Self {
         Self::abs(abs)
     }
 }
 
-impl<T: const UiNum, U: const UiNum> const From<(T, U)> for UiVec2
+const impl<T: const UiNum, U: const UiNum> From<(T, U)> for UiVec2
 where
     (T, U): const Destruct,
 {
@@ -421,7 +421,7 @@ impl Display for UiRegion {
     }
 }
 
-#[derive(Debug)]
+#[derive(Debug, Clone, Copy, PartialEq)]
 pub struct PixelRegion {
     pub top_left: Vec2,
     pub bot_right: Vec2,

core/src/primitive/color.rs

@@ -10,6 +10,15 @@ pub struct Color<T> {
     pub a: T,
 }
 
+/// Required by parley's `Brush`, which every text style is generic over. Opaque
+/// black rather than transparent: a brush that was never set should be visible
+/// and obviously unstyled, not invisible.
+impl<T: ColorNum> Default for Color<T> {
+    fn default() -> Self {
+        Self::BLACK
+    }
+}
+
 impl<T: ColorNum> Color<T> {
     pub const BLACK: Self = Self::rgb(T::MIN, T::MIN, T::MIN);
     pub const WHITE: Self = Self::rgb(T::MAX, T::MAX, T::MAX);
@@ -144,7 +153,7 @@ impl ColorNum for f32 {
 
 unsafe impl bytemuck::Pod for Color<u8> {}
 
-impl const F32Conversion for f32 {
+const impl F32Conversion for f32 {
     fn to(self) -> f32 {
         self
     }
@@ -153,7 +162,7 @@ impl const F32Conversion for f32 {
     }
 }
 
-impl const F32Conversion for u8 {
+const impl F32Conversion for u8 {
     fn to(self) -> f32 {
         self as f32 / 255.0
     }

core/src/primitive/layer.rs

@@ -1,9 +1,6 @@
 use std::ops::{Index, IndexMut};
 
-use crate::{
-    render::{MaskIdx, Primitive, PrimitiveHandle, PrimitiveInst, Primitives},
-    util::to_mut,
-};
+use crate::{render::LayerOrder, util::to_mut};
 
 pub type LayerId = usize;
 
@@ -39,7 +36,10 @@ struct Child {
     tail: usize,
 }
 
-pub type PrimitiveLayers = Layers<Primitives>;
+/// The draw order of every layer. The primitives themselves live in one
+/// arena beside this (`UiRenderState::primitives`); a layer names the
+/// slots it draws, which is what its vertex buffer is.
+pub type PrimitiveLayers = Layers<LayerOrder>;
 
 impl<T: Default> Layers<T> {
     pub fn new() -> Layers<T> {
@@ -119,20 +119,6 @@ impl<T: Default> Layers<T> {
     }
 }
 
-impl PrimitiveLayers {
-    pub fn write<P: Primitive>(
-        &mut self,
-        layer: LayerId,
-        info: PrimitiveInst<P>,
-    ) -> PrimitiveHandle {
-        self[layer].write(layer, info)
-    }
-
-    pub fn free(&mut self, h: &PrimitiveHandle) -> MaskIdx {
-        self[h.layer].free(h)
-    }
-}
-
 impl<T: Default> Default for Layers<T> {
     fn default() -> Self {
         Self::new()

core/src/primitive/text.rs

@@ -1,60 +1,444 @@
-use crate::{Align, RegionAlign, TextureHandle, Textures, UiColor, util::Vec2};
-use cosmic_text::{
-    Attrs, AttrsList, Buffer, CacheKey, Color, Family, FontSystem, Metrics, Placement, SwashCache,
-    SwashContent,
+use crate::{Align, GlyphAtlas, GlyphKey, PlacedGlyph, RegionAlign, Textures, UiColor, util::Vec2};
+use parley::{
+    Alignment, AlignmentOptions, FontContext, FontFamily, FontFamilyName, FontStyle, FontWeight,
+    GenericFamily, Layout, LayoutContext, LineHeight, PositionedLayoutItem, StyleProperty,
+    fontique::Blob,
+};
+use std::ops::Range;
+use std::sync::Arc;
+use swash::{
+    FontRef,
+    scale::{Render, ScaleContext, Source, StrikeWith},
+    zeno::{Format, Vector},
 };
-use image::{DynamicImage, GenericImageView, RgbaImage};
-use std::simd::{Simd, num::SimdUint};
 
-/// TODO: properly wrap this
-pub mod text_lib {
-    pub use cosmic_text::*;
+/// The icon font iris ships: the Nerd Fonts Symbols **Mono** subset built
+/// by `iris/core/build-icon-font.sh`, holding only the codepoints
+/// `crate::icon` names (992 bytes for three glyphs today).
+///
+/// This is the one font bundled here, and it is not a text font: body and
+/// monospace text still come from the platform's own collection
+/// (DECISIONS.md, 2026-09-07). An icon is the opposite case -- a small,
+/// closed set of codepoints no system font is guaranteed to have -- which
+/// is the same division the Compose app makes.
+const NERD_ICONS: &[u8] = include_bytes!("../../assets/fonts/nerd_icons.ttf");
+
+/// What starting up found about text rendering, for the on-screen
+/// Diagnostics page and the one startup log line (RUST.md's P0 box, "log
+/// once at startup ... the number of font families found, the default
+/// family resolved"). Built once by `TextData::font_diagnostics` --
+/// `Default::default` still exists for callers (tests, examples) that
+/// don't need the report.
+#[derive(Clone, Debug)]
+pub struct FontDiagnostics {
+    /// `Collection::family_names().count()` after registering the bundled
+    /// fonts -- system families plus the two bundled ones.
+    pub families_found: usize,
+    /// The family `GenericFamily::SansSerif` resolves to first -- the
+    /// bundled "Noto Sans" unless registration itself failed.
+    pub default_family: Option<String>,
+    /// The family `GenericFamily::Monospace` resolves to first.
+    pub default_mono_family: Option<String>,
+    /// One resolved family name per style axis this crate actually uses
+    /// (`SpanStyle::bold`/`italic`), so a report can say plainly whether a
+    /// bold/italic request is landing on a real face rather than being
+    /// silently absorbed by whatever the sans-serif default resolves to
+    /// for every weight (RUST.md's P0 box, "bold words render as blank
+    /// gaps" -- a family that resolves but has no distinct bold face is
+    /// exactly what produced that).
+    pub regular_resolved: Option<String>,
+    pub bold_resolved: Option<String>,
+    pub italic_resolved: Option<String>,
+    pub mono_resolved: Option<String>,
+    /// The family the bundled icon font registered under, or `None` if
+    /// registering it failed. Reported rather than assumed: it is the one
+    /// font iris ships, so `None` is a broken build and must not look
+    /// like a device that happens to lack a face.
+    pub icon_family: Option<String>,
 }
 
+/// Everything text needs that outlives one string: the font collection, the
+/// layout scratch space, the glyph rasteriser and the atlas they fill.
 pub struct TextData {
-    pub font_system: FontSystem,
-    pub swash_cache: SwashCache,
-    glyph_cache: Vec<(Placement, CacheKey, Color)>,
+    pub font_cx: FontContext,
+    pub layout_cx: LayoutContext<UiColor>,
+    scale_cx: ScaleContext,
+    pub atlas: GlyphAtlas,
+    /// Physical pixels per dp -- a second copy of
+    /// `UiRenderState::density`, kept here too because `TextEditCtx::layout`
+    /// (cursor movement and hit-testing, `widget/text/edit.rs`) shapes text
+    /// from an event callback that has a `TextData` but no `Painter`, so it
+    /// has nowhere else to read the display's density from. Both copies are
+    /// set together, from the one place either backend learns the real
+    /// value (`android::view::new_peer`); this is the same accepted
+    /// duplication as `AndroidRenderer::content_scale`; a single source of
+    /// truth would mean carrying a `Painter` (or output size) into every
+    /// input handler for the sake of one field.
+    pub density: f32,
+    /// The family name [`NERD_ICONS`] registered under, which is what
+    /// [`Family::Icons`] resolves to. `None` only if registering the
+    /// bundled font failed, which is a broken build rather than a
+    /// platform difference -- said in the startup diagnostics rather than
+    /// silently drawn as tofu.
+    pub icon_family: Option<String>,
 }
 
 impl Default for TextData {
+    /// Text comes entirely from the platform's own font collection --
+    /// `FontContext::new()` builds a `fontique::Collection` with
+    /// `CollectionOptions::system_fonts` on by default, which is real
+    /// discovery on both targets this crate ships on: Android's backend
+    /// parses `/system/fonts` and `/system/etc/fonts.xml` and maps
+    /// `SansSerif`/`SystemUi` to `["Roboto Flex", "Roboto", "Noto Sans"]`
+    /// and `Monospace` to the platform's `"monospace"` alias; the desktop
+    /// build's backend is fontconfig. No font is bundled or registered
+    /// here -- see DECISIONS.md's 2026-09-07 entry for why (matching what
+    /// the Compose app does: it takes body/monospace text from
+    /// `FontFamily.Default`/`FontFamily.Monospace`, i.e. Android's Roboto
+    /// and its platform monospace face, and ships no text font of its own,
+    /// only its committed Nerd Fonts icon subset for fixed glyphs).
     fn default() -> Self {
+        let mut font_cx = FontContext::new();
+        patch_android_monospace(&mut font_cx);
+        let icon_family = register_icon_font(&mut font_cx);
         Self {
-            font_system: FontSystem::new(),
-            swash_cache: SwashCache::new(),
-            glyph_cache: Default::default(),
+            font_cx,
+            layout_cx: LayoutContext::new(),
+            scale_cx: ScaleContext::new(),
+            atlas: GlyphAtlas::default(),
+            density: 1.0,
+            icon_family,
         }
     }
 }
 
-#[derive(Clone, Copy)]
+/// Registers the bundled icon font as an ordinary named family and
+/// answers the name it registered under -- read back from the collection
+/// rather than written down here, so the name cannot drift from the file
+/// (`build-icon-font.sh` takes whatever face the Nerd Fonts release
+/// ships).
+///
+/// A *named* family rather than a generic one: nothing should fall back
+/// to it for ordinary text, and nothing should fall back out of it for an
+/// icon -- a system face that happens to have one of these codepoints
+/// would draw somebody else's picture.
+fn register_icon_font(font_cx: &mut FontContext) -> Option<String> {
+    let blob = Blob::new(Arc::new(NERD_ICONS));
+    let id = font_cx
+        .collection
+        .register_fonts(blob, None)
+        .into_iter()
+        .map(|(id, _)| id)
+        .next()?;
+    font_cx.collection.family_name(id).map(str::to_string)
+}
+
+/// Works around `fontique` 0.11.1's Android backend never resolving
+/// `GenericFamily::Monospace` (confirmed against
+/// `fontique-0.11.1/src/backend/android.rs`'s `SystemFonts::new`, and still
+/// present on `linebender/parley`'s `main` as of 2026-09-07, so there is no
+/// released fix to bump to yet -- see DECISIONS.md's 2026-09-07 entry,
+/// "Platform fonts," for the full account). Two bugs stack, not one:
+/// `DEFAULT_GENERIC_FAMILIES` looks up the name `"monospace"` *before*
+/// `fonts.xml` is parsed into that same name map, and even after parsing,
+/// AOSP's `fonts.xml` names it with a `<family name="monospace">` element
+/// (not an `<alias>`) whose `<font>` children the backend's own parser
+/// does not read (a `TODO` in that match arm) -- so the name gets a
+/// `FamilyId` with no font data behind it, and `family_by_name("monospace")`
+/// comes back empty too. Confirmed on this checkout's emulator: `adb pull
+/// /system/etc/fonts.xml` shows
+/// `<family name="monospace"><font weight="400"
+/// style="normal">DroidSansMono.ttf</font></family>` with no matching
+/// alias.
+///
+/// So this reads `fonts.xml` itself (already on-device, already the
+/// authority Compose's own `Typeface.MONOSPACE` resolves through) for the
+/// filename that declaration names, then finds which of fontique's
+/// *actually* scanned families (from `/system/fonts`, which do carry real
+/// font data, just under whatever name the font's own metadata gives it --
+/// "Droid Sans Mono" here, but that name is never hardcoded) owns a font
+/// file with that name, and registers that family as the `Monospace`
+/// generic the way the backend itself would have if its parser had reified
+/// the declaration. A no-op if the family is somehow already resolved
+/// (future fontique) or nothing matches (no `fonts.xml`, e.g. a headless
+/// test, or a device that names it some other way).
+#[cfg(target_os = "android")]
+fn patch_android_monospace(font_cx: &mut FontContext) {
+    use parley::fontique::SourceKind;
+
+    let already_resolved = font_cx
+        .collection
+        .generic_families(GenericFamily::Monospace)
+        .next()
+        .is_some();
+    if already_resolved {
+        return;
+    }
+    let Some(target_file) = android_monospace_font_filename() else {
+        return;
+    };
+    let names: Vec<String> = font_cx
+        .collection
+        .family_names()
+        .map(str::to_string)
+        .collect();
+    for name in names {
+        let Some(id) = font_cx.collection.family_id(&name) else {
+            continue;
+        };
+        let Some(info) = font_cx.collection.family(id) else {
+            continue;
+        };
+        let Some(font) = info.default_font() else {
+            continue;
+        };
+        let SourceKind::Path(path) = font.source().kind() else {
+            continue;
+        };
+        if path.file_name().and_then(|f| f.to_str()) == Some(target_file.as_str()) {
+            font_cx
+                .collection
+                .append_generic_families(GenericFamily::Monospace, std::iter::once(id));
+            return;
+        }
+    }
+}
+
+/// Reads the font filename `fonts.xml` names for its `"monospace"` family
+/// (e.g. `"DroidSansMono.ttf"`), by plain substring search rather than a
+/// real XML parser -- a new dependency for one well-known, stable AOSP file
+/// whose structure fontique itself already parses with a full parser one
+/// module over. Not a general XML reader; assumes the file has exactly one
+/// `<family name="monospace">` element with at least one `<font>` child,
+/// which is the format on every AOSP `fonts.xml` this was checked against.
+#[cfg(target_os = "android")]
+fn android_monospace_font_filename() -> Option<String> {
+    let android_root = std::env::var("ANDROID_ROOT").unwrap_or_else(|_| "/system".to_string());
+    let xml =
+        std::fs::read_to_string(std::path::Path::new(&android_root).join("etc/fonts.xml")).ok()?;
+    let family_start = xml.find("<family name=\"monospace\">")?;
+    let block = &xml[family_start..];
+    let block = &block[..block.find("</family>")?];
+    let font_tag = block.find("<font")?;
+    let after_tag = &block[font_tag..];
+    let content_start = after_tag.find('>')? + 1;
+    let content = &after_tag[content_start..];
+    let filename = content[..content.find('<')?].trim();
+    (!filename.is_empty()).then(|| filename.to_string())
+}
+
+#[cfg(not(target_os = "android"))]
+fn patch_android_monospace(_font_cx: &mut FontContext) {}
+
+impl TextData {
+    /// [`Family::Icons`] as the name the bundled font actually registered
+    /// under; everything else unchanged.
+    ///
+    /// Cloned rather than borrowed because the caller needs it while the
+    /// layout builder holds `&mut self` -- a `String` per shaped icon run,
+    /// paid only when the layout is rebuilt.
+    pub fn resolve_family(&self, family: &Family) -> Family {
+        match family {
+            Family::Icons => self
+                .icon_family
+                .clone()
+                .map_or(Family::Icons, Family::Named),
+            other => other.clone(),
+        }
+    }
+
+    /// Builds the startup report -- see `FontDiagnostics`. Queries the
+    /// collection directly (`fontique::Query`) rather than shaping a real
+    /// string, since all that's needed is which family each axis lands on.
+    pub fn font_diagnostics(&mut self) -> FontDiagnostics {
+        use parley::fontique::{Attributes, FontWidth, QueryStatus};
+        let families_found = self.font_cx.collection.family_names().count();
+        let default_family_id = self
+            .font_cx
+            .collection
+            .generic_families(GenericFamily::SansSerif)
+            .next();
+        let default_family = default_family_id
+            .and_then(|id| self.font_cx.collection.family_name(id).map(str::to_string));
+        let default_mono_family_id = self
+            .font_cx
+            .collection
+            .generic_families(GenericFamily::Monospace)
+            .next();
+        let default_mono_family = default_mono_family_id
+            .and_then(|id| self.font_cx.collection.family_name(id).map(str::to_string));
+
+        // Resolves the family a (generic family, weight, style) query lands
+        // on, without holding the `Query`'s borrow of `collection` across
+        // the `family_name` lookup that needs it back -- the `FamilyId` is
+        // captured out of the closure first, then looked up once `query`
+        // (and its borrow) has been dropped.
+        let mut resolve_family =
+            |generic: GenericFamily, weight: FontWeight, style: FontStyle| -> Option<String> {
+                let mut family_id = None;
+                {
+                    let mut query = self
+                        .font_cx
+                        .collection
+                        .query(&mut self.font_cx.source_cache);
+                    query.set_families([generic]);
+                    query.set_attributes(Attributes {
+                        width: FontWidth::NORMAL,
+                        style,
+                        weight,
+                    });
+                    query.matches_with(|font| {
+                        family_id = Some(font.family.0);
+                        QueryStatus::Stop
+                    });
+                }
+                family_id.and_then(|id| self.font_cx.collection.family_name(id).map(str::to_string))
+            };
+
+        let regular_resolved = resolve_family(
+            GenericFamily::SansSerif,
+            FontWeight::NORMAL,
+            FontStyle::Normal,
+        );
+        let bold_resolved = resolve_family(
+            GenericFamily::SansSerif,
+            FontWeight::BOLD,
+            FontStyle::Normal,
+        );
+        let italic_resolved = resolve_family(
+            GenericFamily::SansSerif,
+            FontWeight::NORMAL,
+            FontStyle::Italic,
+        );
+        let mono_resolved = resolve_family(
+            GenericFamily::Monospace,
+            FontWeight::NORMAL,
+            FontStyle::Normal,
+        );
+
+        FontDiagnostics {
+            families_found,
+            default_family,
+            default_mono_family,
+            regular_resolved,
+            bold_resolved,
+            italic_resolved,
+            mono_resolved,
+            icon_family: self.icon_family.clone(),
+        }
+    }
+}
+
+/// Which family to ask for. Kept as an owned name rather than parley's
+/// borrowed `FontFamily<'_>` so that a widget can hold one without a lifetime.
+#[derive(Clone, PartialEq)]
+pub enum Family {
+    SansSerif,
+    Serif,
+    Monospace,
+    /// The bundled icon font -- see [`crate::icon`] for what is in it.
+    /// Named as an intention rather than as a font name because only
+    /// [`TextData`] knows what the file registered as; it resolves this
+    /// during shaping ([`TextData::resolve_family`]).
+    Icons,
+    Named(String),
+}
+
+impl Family {
+    fn family(&self) -> FontFamily<'_> {
+        let name = match self {
+            Self::SansSerif => FontFamilyName::Generic(GenericFamily::SansSerif),
+            Self::Serif => FontFamilyName::Generic(GenericFamily::Serif),
+            Self::Monospace => FontFamilyName::Generic(GenericFamily::Monospace),
+            // Only reachable if `resolve_family` did not run, which no
+            // shaping path allows -- and sans-serif is the honest answer
+            // for a build whose icon font failed to register: the reader
+            // gets the platform's own tofu rather than a wrong picture.
+            Self::Icons => FontFamilyName::Generic(GenericFamily::SansSerif),
+            Self::Named(name) => FontFamilyName::Named(name.as_str().into()),
+        };
+        FontFamily::Single(name)
+    }
+}
+
+/// One styled run inside a `TextBuffer`, overriding `TextAttrs`' base style
+/// over `range` (a byte range into the buffer's text). Every field is
+/// optional so a span only says what it changes -- e.g. a link span sets
+/// `color` and `underline` and leaves weight/family at the paragraph's own
+/// default. This is I5's answer to RUST.md's inline-rich-text ceiling
+/// (`masonry/src/widgets/text_area.rs`'s `StyleSet` is one style for the
+/// whole editor, with `// TODO: RichTextInput` beside it): parley's own
+/// `RangedBuilder::push` already takes a style and a range, so per-span
+/// bold/italic/monospace/colour/underline only needed plumbing this struct
+/// through to it and giving each glyph its own colour at draw time (see
+/// `PlacedGlyph::color` and `TextData::place` below) instead of the one
+/// `RenderedText::color` every glyph used to share.
+#[derive(Clone, PartialEq)]
+pub struct SpanStyle {
+    pub range: Range<usize>,
+    pub color: Option<UiColor>,
+    pub family: Option<Family>,
+    /// Overrides `TextAttrs::font_size` for just this range -- what lets a
+    /// heading inside a transcript row's single `TextEdit` be bigger than
+    /// the paragraph text around it, so a whole markdown-folded row (block
+    /// and inline styling both) can stay one selectable text buffer instead
+    /// of one widget per block.
+    pub font_size: Option<f32>,
+    pub bold: bool,
+    pub italic: bool,
+    pub underline: bool,
+}
+
+impl SpanStyle {
+    pub fn new(range: Range<usize>) -> Self {
+        Self {
+            range,
+            color: None,
+            family: None,
+            font_size: None,
+            bold: false,
+            italic: false,
+            underline: false,
+        }
+    }
+    pub fn color(mut self, color: UiColor) -> Self {
+        self.color = Some(color);
+        self
+    }
+    pub fn family(mut self, family: Family) -> Self {
+        self.family = Some(family);
+        self
+    }
+    pub fn font_size(mut self, size: f32) -> Self {
+        self.font_size = Some(size);
+        self
+    }
+    pub fn bold(mut self) -> Self {
+        self.bold = true;
+        self
+    }
+    pub fn italic(mut self) -> Self {
+        self.italic = true;
+        self
+    }
+    pub fn underline(mut self) -> Self {
+        self.underline = true;
+        self
+    }
+}
+
+#[derive(Clone, PartialEq)]
 pub struct TextAttrs {
     pub color: UiColor,
     pub font_size: f32,
     pub line_height: f32,
-    pub family: Family<'static>,
+    pub family: Family,
     pub wrap: bool,
     /// inner alignment of text region (within where it's drawn)
     pub align: RegionAlign,
 }
 
-impl TextAttrs {
-    pub fn apply(&self, font_system: &mut FontSystem, buf: &mut Buffer, width: Option<f32>) {
-        buf.set_metrics_and_size(
-            font_system,
-            Metrics::new(self.font_size, self.line_height),
-            width,
-            None,
-        );
-        let attrs = Attrs::new().family(self.family);
-        let list = AttrsList::new(&attrs);
-        for line in &mut buf.lines {
-            line.set_attrs_list(list.clone());
-        }
-    }
-}
-
-pub type TextBuffer = Buffer;
+pub const LINE_HEIGHT_MULT: f32 = 1.1;
 
 impl Default for TextAttrs {
     fn default() -> Self {
@@ -70,122 +454,324 @@ impl Default for TextAttrs {
     }
 }
 
-pub const LINE_HEIGHT_MULT: f32 = 1.1;
+/// A string together with its laid-out form.
+///
+/// The text and the layout live in one place because parley's `Layout` borrows
+/// nothing but is only meaningful against the string it was built from: keeping
+/// them apart is how they get out of step.
+pub struct TextBuffer {
+    text: String,
+    layout: Layout<UiColor>,
+    spans: Vec<SpanStyle>,
+    /// What the current layout was built for, so `shape` can decline to redo
+    /// work that would come out the same. Spans are not part of this key --
+    /// `set_spans` forces `shaped` to `None` directly, the same way `edit`
+    /// does, since spans change far less often than a naive equality check
+    /// on the whole `Vec` would cost to compute every frame.
+    shaped: Option<(TextAttrs, Option<f32>, f32)>,
+}
+
+impl TextBuffer {
+    pub fn new(text: impl Into<String>) -> Self {
+        Self {
+            text: text.into(),
+            layout: Layout::new(),
+            spans: Vec::new(),
+            shaped: None,
+        }
+    }
+
+    /// Replace this buffer's per-range style overrides (I5's rich text --
+    /// see `SpanStyle`). Invalidates the layout unconditionally, mirroring
+    /// `set_text`.
+    pub fn set_spans(&mut self, spans: Vec<SpanStyle>) {
+        self.spans = spans;
+        self.shaped = None;
+    }
+
+    pub fn new_empty() -> Self {
+        Self::new("")
+    }
+
+    pub fn text(&self) -> &str {
+        &self.text
+    }
+
+    pub fn layout(&self) -> &Layout<UiColor> {
+        &self.layout
+    }
+
+    pub fn is_empty(&self) -> bool {
+        self.text.is_empty()
+    }
+
+    pub fn set_text(&mut self, text: impl Into<String>) {
+        let text = text.into();
+        if text != self.text {
+            self.text = text;
+            self.shaped = None;
+        }
+    }
+
+    /// Edit the string in place; invalidates the layout unconditionally, since
+    /// the caller is assumed to have changed something.
+    pub fn edit(&mut self) -> &mut String {
+        self.shaped = None;
+        &mut self.text
+    }
+
+    pub fn size(&self) -> Vec2 {
+        Vec2::new(self.layout.width(), self.layout.height())
+    }
+
+    /// Lay the text out, unless it is already laid out for these
+    /// attributes, this width and this density.
+    ///
+    /// **`attrs.font_size`/`line_height` and every span's own `font_size`
+    /// are density-independent (dp) units, multiplied by `density` here --
+    /// the one place text crosses from the widget tree's dp sizes into the
+    /// physical pixels the shaper and rasteriser (`TextData::place`) both
+    /// then work in.** This is what makes glyphs sharp on a dense display:
+    /// before this existed, `font_size` was already a physical-pixel value
+    /// (RUST.md's P0 box's global-scale stopgap resolved density by
+    /// stretching the whole rendered frame afterward instead), so a glyph
+    /// was rasterised small and then upscaled by whatever the display's
+    /// scale factor was -- exactly the blur Iris's report described.
+    /// Multiplying here instead means the font size hitting `ScaleContext`
+    /// in `place` below is already the display's real physical size, so
+    /// the atlas holds a bitmap at the resolution it is actually shown at.
+    /// `GlyphKey.size` already keys on that resolved `font_size`
+    /// (`(font_size * 16.0).round()`), so a cache entry is naturally per
+    /// physical size with no change needed there.
+    pub fn shape(
+        &mut self,
+        data: &mut TextData,
+        attrs: &TextAttrs,
+        width: Option<f32>,
+        density: f32,
+    ) {
+        if self.shaped.as_ref() == Some(&(attrs.clone(), width, density)) {
+            return;
+        }
+        // Resolved before the builder borrows `data`: `Family::Icons`
+        // names an intention, and the name behind it lives on `TextData`.
+        let base_family = data.resolve_family(&attrs.family);
+        let span_families: Vec<Option<Family>> = self
+            .spans
+            .iter()
+            .map(|span| span.family.as_ref().map(|f| data.resolve_family(f)))
+            .collect();
+        let mut builder = data
+            .layout_cx
+            .ranged_builder(&mut data.font_cx, &self.text, 1.0, true);
+        builder.push_default(StyleProperty::FontFamily(base_family.family()));
+        builder.push_default(StyleProperty::FontSize(attrs.font_size * density));
+        builder.push_default(StyleProperty::LineHeight(LineHeight::Absolute(
+            attrs.line_height * density,
+        )));
+        builder.push_default(StyleProperty::Brush(attrs.color));
+        for (span, family) in self.spans.iter().zip(&span_families) {
+            let range = span.range.clone();
+            if let Some(color) = span.color {
+                builder.push(StyleProperty::Brush(color), range.clone());
+            }
+            if let Some(family) = family {
+                builder.push(StyleProperty::FontFamily(family.family()), range.clone());
+            }
+            if let Some(size) = span.font_size {
+                builder.push(StyleProperty::FontSize(size * density), range.clone());
+            }
+            if span.bold {
+                builder.push(StyleProperty::FontWeight(FontWeight::BOLD), range.clone());
+            }
+            if span.italic {
+                builder.push(StyleProperty::FontStyle(FontStyle::Italic), range.clone());
+            }
+            if span.underline {
+                builder.push(StyleProperty::Underline(true), range.clone());
+            }
+        }
+        builder.build_into(&mut self.layout, &self.text);
+        self.layout.break_all_lines(width);
+        self.layout
+            .align(Alignment::Start, AlignmentOptions::default());
+        self.shaped = Some((attrs.clone(), width, density));
+    }
+}
 
 impl TextData {
-    pub fn draw(
+    /// Rasterise whatever of `buffer` is not in the atlas yet, and return where
+    /// each glyph goes relative to the text's top-left.
+    ///
+    /// Nothing is uploaded for a glyph already in the atlas, which is the point
+    /// of having one: a resize re-runs this and touches the GPU only if the new
+    /// width brought genuinely new glyphs into view.
+    pub fn place(&mut self, buffer: &TextBuffer, textures: &mut Textures) -> Vec<PlacedGlyph> {
+        let mut placed = Vec::new();
+        for line in buffer.layout.lines() {
+            for item in line.items() {
+                let PositionedLayoutItem::GlyphRun(run) = item else {
+                    continue;
+                };
+                let font = run.run().font();
+                let font_size = run.run().font_size();
+                let coords = run.run().normalized_coords();
+                let run_color = run.style().brush;
+                let Some(font_ref) = FontRef::from_index(font.data.as_ref(), font.index as usize)
+                else {
+                    continue;
+                };
+                let coords_hash = hash_coords(coords);
+                // `font.data.id()` rather than the pointer, so the same font
+                // loaded twice is still one set of entries.
+                let font_id = font.data.id();
+
+                for glyph in run.positioned_glyphs() {
+                    let subpixel = ((glyph.x.fract() * 4.0).round() as i32).rem_euclid(4) as u8;
+                    let key = GlyphKey {
+                        font: font_id,
+                        glyph: glyph.id,
+                        size: (font_size * 16.0).round() as u32,
+                        subpixel,
+                        coords: coords_hash,
+                    };
+                    let entry = match self.atlas.get(&key) {
+                        Some(entry) => entry,
+                        None => {
+                            let mut scaler = self
+                                .scale_cx
+                                .builder(font_ref)
+                                .size(font_size)
+                                .hint(true)
+                                .normalized_coords(coords)
+                                .build();
+                            let image = Render::new(&[
+                                Source::ColorOutline(0),
+                                Source::ColorBitmap(StrikeWith::BestFit),
+                                Source::Outline,
+                            ])
+                            .format(Format::Alpha)
+                            .offset(Vector::new(subpixel as f32 / 4.0, 0.0))
+                            .render(&mut scaler, glyph.id as u16);
+                            match image {
+                                Some(image) => self.atlas.insert(key, &image, textures),
+                                None => {
+                                    self.atlas.insert_empty(key);
+                                    None
+                                }
+                            }
+                        }
+                    };
+                    let Some(entry) = entry else { continue };
+                    placed.push(PlacedGlyph {
+                        entry,
+                        offset: Vec2::new(
+                            glyph.x.floor() + entry.left as f32,
+                            glyph.y.floor() - entry.top as f32,
+                        ),
+                        color: run_color,
+                    });
+                }
+            }
+        }
+        placed
+    }
+}
+
+fn hash_coords(coords: &[i16]) -> u64 {
+    // FxHash over the coordinates; they are short and change rarely.
+    let mut h: u64 = 0xcbf2_9ce4_8422_2325;
+    for c in coords {
+        h ^= *c as u16 as u64;
+        h = h.wrapping_mul(0x1000_0000_01b3);
+    }
+    h
+}
+
+/// A laid-out string, ready to draw: where each glyph goes and how big the
+/// whole thing is.
+///
+/// Cheap to clone and to keep, which is the point -- a widget holds one across
+/// frames and re-emits its quads without going near the rasteriser. `color`
+/// is the buffer's *base* colour (`TextAttrs::color`) for a caller that wants
+/// it as a whole (e.g. tinting a cursor to match); the colour each glyph is
+/// actually drawn in is `PlacedGlyph::color`, which a `SpanStyle` can
+/// override per range.
+#[derive(Clone)]
+pub struct RenderedText {
+    pub glyphs: std::sync::Arc<Vec<PlacedGlyph>>,
+    pub size: Vec2,
+    pub color: UiColor,
+    /// The [`GlyphAtlas::generation`] the glyphs above were placed against.
+    /// A holder must re-render rather than re-emit these quads once the
+    /// atlas has moved on (`GlyphAtlas::clear`'s doc says what happens
+    /// otherwise); `Painter::glyphs` debug-asserts it.
+    pub generation: u64,
+}
+
+impl TextData {
+    /// Lay out and place in one step, which is what a widget wants.
+    pub fn render(
         &mut self,
         buffer: &mut TextBuffer,
         attrs: &TextAttrs,
+        width: Option<f32>,
         textures: &mut Textures,
+        density: f32,
     ) -> RenderedText {
-        // TODO: either this or the layout stuff (or both) is super slow,
-        // should probably do texture packing and things if possible.
-        // very visible if you add just a couple of wrapping texts and resize window
-        // should also be timed to figure out exactly what points need to be sped up
-        // let mut pixels = HashMap::<_, [u8; 4]>::default();
-        let mut min_x = 0;
-        let mut min_y = 0;
-        let mut max_x = 0;
-        let mut max_y = 0;
-        let text_color = {
-            let c = attrs.color;
-            cosmic_text::Color::rgba(c.r, c.g, c.b, c.a)
-        };
-        let mut max_width = 0.0f32;
-        let mut height = 0.0;
-
-        for run in buffer.layout_runs() {
-            for glyph in run.glyphs.iter() {
-                let physical_glyph = glyph.physical((0., 0.), 1.0);
-
-                let glyph_color = match glyph.color_opt {
-                    Some(some) => some,
-                    None => text_color,
-                };
-
-                if let Some(img) = self
-                    .swash_cache
-                    .get_image(&mut self.font_system, physical_glyph.cache_key)
-                {
-                    let mut pos = img.placement;
-                    pos.left += physical_glyph.x;
-                    pos.top = physical_glyph.y + run.line_y as i32 - pos.top;
-                    min_x = min_x.min(pos.left);
-                    min_y = min_y.min(pos.top);
-                    max_x = max_x.max(pos.left + pos.width as i32);
-                    max_y = max_y.max(pos.top + pos.height as i32);
-                    self.glyph_cache
-                        .push((pos, physical_glyph.cache_key, glyph_color));
-                }
-            }
-            max_width = max_width.max(run.line_w);
-            height += run.line_height;
-        }
-        let img_width = (max_x - min_x + 1) as u32;
-        let img_height = (max_y - min_y + 1) as u32;
-        let mut image = RgbaImage::new(img_width, img_height);
-
-        for (pos, key, color) in self.glyph_cache.drain(..) {
-            let img = self
-                .swash_cache
-                .get_image(&mut self.font_system, key)
-                .as_ref()
-                .unwrap();
-            let mut merge = |i, color: [u8; 4]| {
-                let i = i as i32;
-                let x = (i % pos.width as i32 + pos.left - min_x) as u32;
-                let y = (i / pos.width as i32 + pos.top - min_y) as u32;
-                let pixel = &mut image[(x, y)].0;
-                // TODO: no clue if proper alpha blending should be done
-                *pixel = Simd::from(color).saturating_add(Simd::from(*pixel)).into();
-            };
-
-            match img.content {
-                SwashContent::Mask => {
-                    for (i, a) in img.data.iter().enumerate() {
-                        let mut color = color.as_rgba();
-                        color[3] = ((color[3] as u32 * *a as u32) / u8::MAX as u32) as u8;
-                        merge(i, color);
-                    }
-                }
-                SwashContent::SubpixelMask => todo!("subpixel mask text rendering"),
-                SwashContent::Color => {
-                    let (colors, _) = img.data.as_chunks::<4>();
-                    for (i, color) in colors.iter().enumerate() {
-                        merge(i, *color);
-                    }
-                }
-            }
-        }
-
-        let max_dim = 8192;
-        if image.width() > max_dim || image.height() > max_dim {
-            let width = image.width().min(max_dim);
-            let height = image.height().min(max_dim);
-            eprintln!(
-                "WARNING: image of size {:?} cropped to {:?} (texture too big)",
-                image.dimensions(),
-                (width, height)
-            );
-            image = image.view(0, 0, width, height).to_image();
-        }
-
+        buffer.shape(self, attrs, width, density);
+        let glyphs = self.place(buffer, textures);
         RenderedText {
-            handle: textures.add(image),
-            top_left_offset: Vec2::new(min_x as f32, min_y as f32),
-            size: Vec2::new(max_width, height),
+            glyphs: std::sync::Arc::new(glyphs),
+            size: buffer.size(),
+            color: attrs.color,
+            generation: self.atlas.generation(),
         }
     }
 }
 
-#[derive(Clone)]
-pub struct RenderedText {
-    pub handle: TextureHandle,
-    pub top_left_offset: Vec2,
-    pub size: Vec2,
-}
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::icon;
 
-pub trait HasTextures {
-    fn add_texture(&mut self, image: DynamicImage) -> TextureHandle;
+    /// Every codepoint `icon` names is actually in the subset the script
+    /// built. This is the failure `build-icon-font.sh`'s own comment warns
+    /// about -- a constant added on one side and not the other is a glyph
+    /// that silently isn't there -- and it is invisible at runtime,
+    /// because a missing glyph draws as nothing rather than as an error.
+    #[test]
+    fn every_icon_is_in_the_bundled_font() {
+        let font = FontRef::from_index(NERD_ICONS, 0).expect("the bundled icon font parses");
+        let charmap = font.charmap();
+        for (name, glyph) in [
+            ("OPEN", icon::OPEN),
+            ("CLOSED", icon::CLOSED),
+            ("COLLAPSE", icon::COLLAPSE),
+        ] {
+            let mut chars = glyph.chars();
+            let ch = chars.next().expect("an icon is one character");
+            assert!(chars.next().is_none(), "{name} is more than one character");
+            assert_ne!(
+                charmap.map(ch),
+                0,
+                "{name} (U+{:04X}) is not in nerd_icons.ttf -- add it to \
+                 build-icon-font.sh's GLYPHS and rerun the script",
+                ch as u32
+            );
+        }
+    }
+
+    /// The font registers, so `Family::Icons` resolves to a real family
+    /// rather than falling through to sans-serif and drawing tofu.
+    #[test]
+    fn the_icon_family_registers_and_resolves() {
+        let data = TextData::default();
+        let family = data.resolve_family(&Family::Icons);
+        assert!(
+            matches!(family, Family::Named(_)),
+            "the bundled icon font did not register: {:?}",
+            data.icon_family
+        );
+    }
 }

core/src/primitive/texture.rs

@@ -1,19 +1,44 @@
-use crate::{
-    render::TexturePrimitive,
-    util::{RefCounter, Vec2},
-};
+use crate::util::{RefCounter, Vec2};
 use image::{DynamicImage, GenericImageView};
 use std::{
+    collections::HashMap,
     ops::Index,
     sync::mpsc::{Receiver, Sender, channel},
 };
 
+/// Which of the two things a texture slot holds. See TEXTURES.md's
+/// "Recommended shape" for why these are drawn so differently: a page is a
+/// layer of one shared array texture and never gets its own bind group; a
+/// standalone image is the opposite, one texture and one bind group, never a
+/// layer.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub enum TextureKind {
+    Image,
+    /// The array-texture layer this page was assigned. Chosen synchronously
+    /// by `Textures::add_page` rather than by the renderer, because glyph
+    /// insertion needs it in the same call, before any GPU sync happens.
+    Page {
+        layer: u32,
+    },
+}
+
+/// What a [`Textures::shared`] texture is a picture of -- exactly, not by
+/// hash: `owner` names the widget kind whose description it is, and `id`
+/// packs that description's own fields, so two owners cannot collide and
+/// a debugger shows which picture a slot holds.
+#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
+pub struct SharedTextureKey {
+    pub owner: &'static str,
+    pub id: u64,
+}
+
 #[derive(Debug, Clone)]
 pub struct TextureHandle {
-    inner: TexturePrimitive,
+    slot: u32,
+    kind: TextureKind,
     size: Vec2,
     counter: RefCounter,
-    send: Sender<u32>,
+    send: Sender<(TextureKind, u32)>,
 }
 
 /// a texture manager for a ui
@@ -21,22 +46,47 @@ pub struct TextureHandle {
 pub struct Textures {
     free: Vec<u32>,
     images: Vec<Option<DynamicImage>>,
+    /// What each slot is, kept beside the image so a slot can be pushed
+    /// again without the handle that knows -- see [`Textures::reupload`].
+    kinds: Vec<TextureKind>,
+    /// Textures built from a description rather than from a file, one per
+    /// distinct description: see [`Textures::shared`]. The map holds a
+    /// reference of its own, so a shared texture outlives every widget
+    /// drawing it and its slot is never recycled underneath one.
+    shared: HashMap<SharedTextureKey, TextureHandle>,
+    /// Next layer to hand out to an atlas page. Pages are never freed (no
+    /// atlas eviction), so this only grows and `free` never holds one.
+    next_page_layer: u32,
     updates: Vec<Update>,
-    send: Sender<u32>,
-    recv: Receiver<u32>,
+    send: Sender<(TextureKind, u32)>,
+    recv: Receiver<(TextureKind, u32)>,
 }
 
 pub enum TextureUpdate<'a> {
-    Push(&'a DynamicImage),
-    Set(u32, &'a DynamicImage),
+    Push(TextureKind, &'a DynamicImage),
+    Set(TextureKind, u32, &'a DynamicImage),
+    /// Overwrite a rectangle of an existing texture, rather than replacing it.
+    /// The glyph atlas grows a glyph at a time, and re-uploading a whole atlas
+    /// per glyph is megabytes of copy for a few hundred bytes of change.
+    /// Only ever issued against a page -- a standalone image is never patched.
+    Patch(u32, PatchRect, &'a DynamicImage),
     Free(u32),
-    PushFree,
+    PushFree(TextureKind),
     SetFree,
 }
 
+#[derive(Debug, Clone, Copy)]
+pub struct PatchRect {
+    pub x: u32,
+    pub y: u32,
+    pub width: u32,
+    pub height: u32,
+}
+
 enum Update {
-    Push(u32),
-    Set(u32),
+    Push(TextureKind, u32),
+    Set(TextureKind, u32),
+    Patch(u32, PatchRect),
     Free(u32),
 }
 
@@ -46,58 +96,162 @@ impl Textures {
         Self {
             free: Vec::new(),
             images: Vec::new(),
+            kinds: Vec::new(),
+            shared: HashMap::new(),
+            next_page_layer: 0,
             updates: Vec::new(),
             send,
             recv,
         }
     }
+
     pub fn add(&mut self, image: impl Into<DynamicImage>) -> TextureHandle {
         let image = image.into();
         let size = image.dimensions().into();
-        let view_idx = self.push(image);
-        // 0 == default in renderer; TODO: actually create samplers here
-        let sampler_idx = 0;
+        let kind = TextureKind::Image;
+        let slot = self.push(kind, image);
         TextureHandle {
-            inner: TexturePrimitive {
-                view_idx,
-                sampler_idx,
-            },
+            slot,
+            kind,
             size,
             counter: RefCounter::new(),
             send: self.send.clone(),
         }
     }
 
-    fn push(&mut self, image: DynamicImage) -> u32 {
+    /// Adds a page of the shared glyph atlas array. Only `atlas.rs` should
+    /// call this -- everything else wants `add`.
+    pub fn add_page(&mut self, image: impl Into<DynamicImage>) -> TextureHandle {
+        let image = image.into();
+        let size = image.dimensions().into();
+        let layer = self.next_page_layer;
+        self.next_page_layer += 1;
+        let kind = TextureKind::Page { layer };
+        let slot = self.push(kind, image);
+        TextureHandle {
+            slot,
+            kind,
+            size,
+            counter: RefCounter::new(),
+            send: self.send.clone(),
+        }
+    }
+
+    fn push(&mut self, kind: TextureKind, image: DynamicImage) -> u32 {
         if let Some(i) = self.free.pop() {
             self.images[i as usize] = Some(image);
-            self.updates.push(Update::Set(i));
+            self.kinds[i as usize] = kind;
+            self.updates.push(Update::Set(kind, i));
             i
         } else {
             let i = self.images.len() as u32;
             self.images.push(Some(image));
-            self.updates.push(Update::Push(i));
+            self.kinds.push(kind);
+            self.updates.push(Update::Push(kind, i));
             i
         }
     }
 
+    /// The one texture for `key`, building it on the first ask and handing
+    /// out a further reference to it every time after.
+    ///
+    /// **Why this exists**: a texture rasterised from a *description* --
+    /// `widget::mark`'s triangle, from a direction and a colour -- has as
+    /// many copies as there are widgets asking for it, and each copy is
+    /// its own GPU texture, its own bind group and its own draw call. A
+    /// transcript screen with a folded card per tool call built one per
+    /// card: hundreds of 48x48 textures of three distinct pictures,
+    /// created and freed again as rows recycled. `make` is not called when
+    /// the key is already known, so the rasterising is paid once too.
+    ///
+    /// The map keeps its own reference for the life of the `Textures`, so
+    /// a shared slot is never freed and never reused for something else --
+    /// which is what makes a handle held by a long-lived widget safe.
+    pub fn shared(
+        &mut self,
+        key: SharedTextureKey,
+        make: impl FnOnce() -> DynamicImage,
+    ) -> TextureHandle {
+        if let Some(handle) = self.shared.get(&key) {
+            return handle.clone();
+        }
+        let handle = self.add(make());
+        self.shared.insert(key, handle.clone());
+        handle
+    }
+
+    /// The stored image for a handle, to be written into before `patch`.
+    pub fn image_mut(&mut self, handle: &TextureHandle) -> &mut DynamicImage {
+        self.images[handle.slot as usize]
+            .as_mut()
+            .expect("texture was freed while still held")
+    }
+
+    /// Queue an upload of just `rect`, after writing it with `image_mut`.
+    pub fn patch(&mut self, handle: &TextureHandle, rect: PatchRect) {
+        self.updates.push(Update::Patch(handle.slot, rect));
+    }
+
+    /// Queue every live slot for upload again, in slot order -- what a
+    /// genuinely new GPU device needs, in place of forgetting everything.
+    ///
+    /// A new device starts with no textures, and the renderer-side mirror
+    /// of these slots (`render::texture::GpuTextures`) starts empty with
+    /// it. What it must not do is start empty while the handles widgets
+    /// are still holding name slots by *index*: `Textures::reset` used to
+    /// throw this bookkeeping away, which left every live `TextureHandle`
+    /// -- one per `widget::mark`, hundreds on a transcript screen --
+    /// pointing at a slot nothing recognised, and the first frame after an
+    /// Android surface rebuild panicked in `image_bind_group` ("texture
+    /// slot 89 is not a live standalone image: None"). Re-uploading
+    /// instead keeps every index meaning what it meant, because this side
+    /// still holds the images: the slot list is rebuilt identically,
+    /// including the empty slots, which go across as `PushFree` so the
+    /// ones after them still land where they were.
+    ///
+    /// The glyph atlas comes back with it and is deliberately *not*
+    /// cleared any more: its pages are slots here, this side holds their
+    /// pixels, and re-uploading them restores exactly the atlas that was
+    /// there -- so an app switch no longer costs a re-rasterisation of
+    /// every glyph on screen either.
+    ///
+    /// Pending updates are dropped rather than kept: each is either a push
+    /// or a patch of a slot this replays in full.
+    pub fn reupload(&mut self) {
+        self.updates.clear();
+        self.updates
+            .extend((0..self.images.len() as u32).map(|i| Update::Push(self.kinds[i as usize], i)));
+    }
+
     pub fn free(&mut self) {
-        for idx in self.recv.try_iter() {
+        for (kind, idx) in self.recv.try_iter() {
             self.images[idx as usize] = None;
             self.updates.push(Update::Free(idx));
-            self.free.push(idx);
+            // A page's slot is never reclaimed: `GlyphAtlas` never drops the
+            // handles it holds, and there is no eviction path for a hole in
+            // the middle of the array's layers. If that ever changes, this
+            // is where a freed page's layer would need to go on a free list
+            // of its own, separate from `free`, which only ever holds
+            // ordinary image slots today.
+            if kind == TextureKind::Image {
+                self.free.push(idx);
+            }
         }
     }
 
     pub fn updates(&mut self) -> impl Iterator<Item = TextureUpdate<'_>> {
         self.updates.drain(..).map(|u| match u {
-            Update::Push(i) => self.images[i as usize]
+            Update::Push(kind, i) => self.images[i as usize]
                 .as_ref()
-                .map(TextureUpdate::Push)
-                .unwrap_or(TextureUpdate::PushFree),
-            Update::Set(i) => self.images[i as usize]
+                .map(|img| TextureUpdate::Push(kind, img))
+                .unwrap_or(TextureUpdate::PushFree(kind)),
+            Update::Set(kind, i) => self.images[i as usize]
                 .as_ref()
-                .map(|img| TextureUpdate::Set(i, img))
+                .map(|img| TextureUpdate::Set(kind, i, img))
+                .unwrap_or(TextureUpdate::SetFree),
+            Update::Patch(i, rect) => self.images[i as usize]
+                .as_ref()
+                .map(|img| TextureUpdate::Patch(i, rect, img))
                 .unwrap_or(TextureUpdate::SetFree),
             Update::Free(i) => TextureUpdate::Free(i),
         })
@@ -105,18 +259,36 @@ impl Textures {
 }
 
 impl TextureHandle {
-    pub fn primitive(&self) -> TexturePrimitive {
-        self.inner
-    }
     pub fn size(&self) -> Vec2 {
         self.size
     }
+
+    /// The bind-group index this handle draws with. Only valid for a
+    /// standalone image; an atlas page has no bind group of its own -- it
+    /// samples the shared array via `layer()` instead. Getting this wrong is
+    /// a caller bug (the wrong kind of handle reached the wrong draw path),
+    /// not a recoverable condition, so it panics rather than drawing garbage.
+    pub fn image_index(&self) -> u32 {
+        match self.kind {
+            TextureKind::Image => self.slot,
+            TextureKind::Page { .. } => panic!("image_index() called on an atlas page handle"),
+        }
+    }
+
+    /// The layer this page occupies in the shared atlas array texture.
+    /// Only valid for a page handle; see `image_index`'s note.
+    pub fn layer(&self) -> u32 {
+        match self.kind {
+            TextureKind::Page { layer } => layer,
+            TextureKind::Image => panic!("layer() called on a standalone image handle"),
+        }
+    }
 }
 
 impl Drop for TextureHandle {
     fn drop(&mut self) {
         if self.counter.drop() {
-            let _ = self.send.send(self.inner.view_idx);
+            let _ = self.send.send((self.kind, self.slot));
         }
     }
 }
@@ -125,7 +297,7 @@ impl Index<&TextureHandle> for Textures {
     type Output = DynamicImage;
 
     fn index(&self, index: &TextureHandle) -> &Self::Output {
-        self.images[index.inner.view_idx as usize].as_ref().unwrap()
+        self.images[index.slot as usize].as_ref().unwrap()
     }
 }
 
@@ -134,3 +306,90 @@ impl Default for Textures {
         Self::new()
     }
 }
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use image::RgbaImage;
+
+    fn image(n: u32) -> DynamicImage {
+        RgbaImage::new(n, n).into()
+    }
+
+    fn key(id: u64) -> SharedTextureKey {
+        SharedTextureKey { owner: "test", id }
+    }
+
+    /// What `widget::mark` needs: one texture per description, however
+    /// many widgets ask for it, and a different description is a
+    /// different texture.
+    #[test]
+    fn a_shared_texture_is_built_once_and_handed_out_again() {
+        let mut textures = Textures::new();
+        let built = std::cell::Cell::new(0);
+        let make = |textures: &mut Textures, id: u64| {
+            textures.shared(key(id), || {
+                built.set(built.get() + 1);
+                image(4)
+            })
+        };
+        let first = make(&mut textures, 1);
+        let again = make(&mut textures, 1);
+        let other = make(&mut textures, 2);
+        assert_eq!(built.get(), 2, "the second ask for key 1 rasterised again");
+        assert_eq!(first.image_index(), again.image_index());
+        assert_ne!(first.image_index(), other.image_index());
+    }
+
+    /// The map's own reference is what keeps a shared slot alive: every
+    /// widget holding one can go away and the slot must not be recycled,
+    /// because the next widget to ask gets that same index back.
+    #[test]
+    fn a_shared_slot_is_not_freed_when_the_last_widget_drops_it() {
+        let mut textures = Textures::new();
+        let slot = textures.shared(key(1), || image(4)).image_index();
+        textures.free();
+        let plain = textures.add(image(4));
+        assert_ne!(
+            plain.image_index(),
+            slot,
+            "an ordinary texture was handed the shared mark's slot"
+        );
+    }
+
+    /// A new GPU device gets the same slot numbering back, so a handle a
+    /// widget has been holding all along still names its own texture --
+    /// the crash `reupload` replaced `reset` to fix.
+    #[test]
+    fn reupload_replays_every_slot_in_order_including_the_empty_ones() {
+        let mut textures = Textures::new();
+        let keep_a = textures.add(image(4));
+        let dropped = textures.add(image(4));
+        let keep_b = textures.add(image(4));
+        let (a, gone, b) = (
+            keep_a.image_index(),
+            dropped.image_index(),
+            keep_b.image_index(),
+        );
+        drop(dropped);
+        textures.free();
+        // Drain the updates so far, the way a frame does.
+        assert!(textures.updates().count() > 0);
+
+        textures.reupload();
+        let kinds: Vec<String> = textures
+            .updates()
+            .map(|u| match u {
+                TextureUpdate::Push(..) => "push".to_string(),
+                TextureUpdate::PushFree(..) => "push-free".to_string(),
+                _ => "other".to_string(),
+            })
+            .collect();
+        assert_eq!(
+            kinds,
+            ["push", "push-free", "push"],
+            "slots {a}, {gone} (freed) and {b} must replay in order, so the \
+             indices after a hole still land where they were"
+        );
+    }
+}

core/src/render/atlas.rs

@@ -0,0 +1,284 @@
+//! A glyph atlas: one texture holding many rasterised glyphs, so drawing text
+//! is a quad per glyph rather than a texture per string.
+//!
+//! What this replaces is why it exists. Text used to be rasterised into its own
+//! `RgbaImage` and uploaded as a whole texture, per text widget, every time
+//! anything about it changed -- so every window resize re-rasterised and
+//! re-uploaded every visible string, which is what the TODO meant by "resizing
+//! (per frame) is really slow". Here a glyph is rasterised once for a given
+//! font, size and subpixel offset and then reused by every string that contains
+//! it, and a resize re-emits quads without touching the GPU's copy at all.
+
+use crate::{
+    PatchRect, TextureHandle, Textures, UiColor,
+    util::{HashMap, Vec2},
+};
+use image::RgbaImage;
+use swash::scale::image::{Content, Image};
+
+/// Side of one atlas page, in pixels. 1024 is 4 MB at RGBA8 -- enough for a
+/// few thousand glyphs at UI sizes, and small enough that a page nobody fills
+/// is not a big waste. Also the fixed width/height of every layer of the
+/// shared array texture in `render::texture` -- `pub(crate)` so that module
+/// can size it without a second constant to keep in sync.
+pub(crate) const PAGE: u32 = 1024;
+
+/// Transparent margin kept around every glyph, so that sampling one cannot
+/// pick up its neighbour along a shared edge.
+const PAD: u32 = 1;
+
+/// Identifies a rasterised glyph. Anything that changes the pixels has to be in
+/// here, or two different glyphs share one entry and the wrong one is drawn.
+#[derive(Clone, Copy, PartialEq, Eq, Hash)]
+pub struct GlyphKey {
+    pub font: u64,
+    pub glyph: u32,
+    /// Font size in 1/16 px, so sizes that round to the same pixels share a
+    /// raster instead of filling the atlas with near-duplicates.
+    pub size: u32,
+    /// Horizontal subpixel phase, in 1/4 px.
+    pub subpixel: u8,
+    /// Hash of the variation coordinates; a variable font at two weights is two
+    /// different sets of pixels from one glyph id.
+    pub coords: u64,
+}
+
+#[derive(Clone, Copy)]
+pub struct GlyphEntry {
+    pub uv_min: [f32; 2],
+    pub uv_max: [f32; 2],
+    /// Offset from the glyph's pen position to the top-left of its pixels.
+    pub left: i32,
+    pub top: i32,
+    pub width: u32,
+    pub height: u32,
+    pub is_color: bool,
+    /// The atlas array layer this glyph's page occupies.
+    pub layer: u32,
+}
+
+struct Page {
+    handle: TextureHandle,
+    /// Shelf packing: glyphs are placed left to right along a shelf whose
+    /// height is the tallest glyph on it, and a new shelf starts above when the
+    /// row runs out. Chosen over a real packer because glyphs at one size are
+    /// close to the same height, which is the case shelves are good at.
+    x: u32,
+    y: u32,
+    shelf_height: u32,
+}
+
+#[derive(Default)]
+pub struct GlyphAtlas {
+    pages: Vec<Page>,
+    /// Bumped by [`GlyphAtlas::clear`], so anything holding placed glyphs
+    /// from an earlier atlas can tell that its coordinates are stale --
+    /// see that method's doc for what goes wrong without it.
+    generation: u64,
+    /// `None` for a glyph that rasterised to nothing -- a space, say. Cached
+    /// too, so it is not re-rasterised on every layout.
+    entries: HashMap<GlyphKey, Option<GlyphEntry>>,
+}
+
+impl GlyphAtlas {
+    pub fn get(&self, key: &GlyphKey) -> Option<Option<GlyphEntry>> {
+        self.entries.get(key).copied()
+    }
+
+    /// Rasterised pixels in, a place in the atlas out. `None` means the glyph
+    /// has no pixels, which is a normal answer rather than a failure.
+    pub fn insert(
+        &mut self,
+        key: GlyphKey,
+        image: &Image,
+        textures: &mut Textures,
+    ) -> Option<GlyphEntry> {
+        let w = image.placement.width;
+        let h = image.placement.height;
+        if w == 0 || h == 0 {
+            self.entries.insert(key, None);
+            return None;
+        }
+        if w + PAD * 2 > PAGE || h + PAD * 2 > PAGE {
+            // A single glyph larger than a page. Refusing is better than
+            // silently drawing a cropped one; the caller draws nothing.
+            self.entries.insert(key, None);
+            return None;
+        }
+
+        let (page_idx, x, y) = self.allocate(w, h, textures);
+        let page = &self.pages[page_idx];
+
+        let img = textures.image_mut(&page.handle);
+        let rgba = img.as_mut_rgba8().expect("atlas page is rgba8");
+        write_glyph(rgba, image, x, y);
+
+        let handle = page.handle.clone();
+        let rect = PatchRect {
+            x,
+            y,
+            width: w,
+            height: h,
+        };
+        textures.patch(&handle, rect);
+
+        let page = &self.pages[page_idx];
+        let scale = 1.0 / PAGE as f32;
+        let entry = GlyphEntry {
+            uv_min: [x as f32 * scale, y as f32 * scale],
+            uv_max: [(x + w) as f32 * scale, (y + h) as f32 * scale],
+            left: image.placement.left,
+            top: image.placement.top,
+            width: w,
+            height: h,
+            is_color: matches!(image.content, Content::Color),
+            layer: page.handle.layer(),
+        };
+        self.entries.insert(key, Some(entry));
+        Some(entry)
+    }
+
+    /// A free `w`x`h` spot, opening a shelf or a page as needed.
+    fn allocate(&mut self, w: u32, h: u32, textures: &mut Textures) -> (usize, u32, u32) {
+        let need_w = w + PAD;
+        let need_h = h + PAD;
+        if let Some(i) = self.pages.iter().position(|p| fits(p, need_w, need_h)) {
+            let page = &mut self.pages[i];
+            if page.x + need_w > PAGE {
+                page.y += page.shelf_height;
+                page.x = PAD;
+                page.shelf_height = 0;
+            }
+            let (x, y) = (page.x, page.y);
+            page.x += need_w;
+            page.shelf_height = page.shelf_height.max(need_h);
+            return (i, x, y);
+        }
+
+        let handle = textures.add_page(RgbaImage::new(PAGE, PAGE));
+        self.pages.push(Page {
+            handle,
+            x: PAD + w + PAD,
+            y: PAD,
+            shelf_height: h + PAD,
+        });
+        (self.pages.len() - 1, PAD, PAD)
+    }
+
+    /// Record that a glyph has no pixels, so it is not re-rasterised.
+    pub fn insert_empty(&mut self, key: GlyphKey) {
+        self.entries.insert(key, None);
+    }
+
+    /// Which atlas the entries handed out right now belong to. A
+    /// [`crate::RenderedText`] records this when it is built and is only
+    /// reusable while it still matches.
+    pub fn generation(&self) -> u64 {
+        self.generation
+    }
+
+    pub fn page_count(&self) -> usize {
+        self.pages.len()
+    }
+
+    pub fn glyph_count(&self) -> usize {
+        self.entries.len()
+    }
+
+    /// Forget every page and every rasterised entry -- what a genuinely new
+    /// GPU device needs (`android::view::IrisViewPeer::surface_changed`'s
+    /// "not already live" branch, e.g. after backgrounding): the pages this
+    /// atlas remembers are `TextureHandle`s into the *old* device's
+    /// textures, which no longer exist, and every `GlyphEntry`'s `uv_min`/
+    /// `uv_max`/`layer` point into them. Without this, a glyph already
+    /// cached here is treated as "already placed" and never re-inserted
+    /// into the fresh (empty) atlas the new renderer actually has --
+    /// exactly the "rectangles stay, glyphs disappear" bug the resize path
+    /// (`AndroidRenderer::resize`) was built to avoid for the reuse case;
+    /// this is its counterpart for the case where the renderer really is
+    /// new. Dropping `pages` also drops its `TextureHandle`s, which send a
+    /// free message back through their `Textures`; see `Textures::reset`'s
+    /// doc for why that is harmless here.
+    /// Bumping `generation` here is the other half of the same
+    /// invalidation: emptying this atlas does nothing about the
+    /// `RenderedText`s widgets are *already holding*
+    /// (`iris::widget::TextView`'s `tex` cache), whose `PlacedGlyph`s carry
+    /// `uv_min`/`uv_max`/`layer` into the atlas that has just been thrown
+    /// away. Those redraw perfectly happily and sample whatever now sits at
+    /// those coordinates -- the fragments-of-other-glyphs Iris photographed
+    /// after resuming the app on 2026-09-06. One counter, checked where the
+    /// cache is read, is what makes a cached render un-reusable across a
+    /// renderer rebuild.
+    pub fn clear(&mut self) {
+        self.pages.clear();
+        self.entries.clear();
+        self.generation += 1;
+    }
+}
+
+fn fits(page: &Page, need_w: u32, need_h: u32) -> bool {
+    // On the current shelf, or on a new one above it.
+    (page.x + need_w <= PAGE && page.y + need_h <= PAGE)
+        || (need_w + PAD <= PAGE && page.y + page.shelf_height + need_h <= PAGE)
+}
+
+/// Copy one rasterised glyph into the page image at `(x, y)`.
+///
+/// A mask glyph keeps its coverage in alpha with the colour left to the shader,
+/// so one raster serves text of any colour; a colour glyph carries its own.
+fn write_glyph(page: &mut RgbaImage, image: &Image, x: u32, y: u32) {
+    let w = image.placement.width;
+    let h = image.placement.height;
+    match image.content {
+        Content::Mask => {
+            for row in 0..h {
+                for col in 0..w {
+                    let a = image.data[(row * w + col) as usize];
+                    page.put_pixel(x + col, y + row, image::Rgba([255, 255, 255, a]));
+                }
+            }
+        }
+        Content::Color => {
+            for row in 0..h {
+                for col in 0..w {
+                    let i = ((row * w + col) * 4) as usize;
+                    let px = [
+                        image.data[i],
+                        image.data[i + 1],
+                        image.data[i + 2],
+                        image.data[i + 3],
+                    ];
+                    page.put_pixel(x + col, y + row, image::Rgba(px));
+                }
+            }
+        }
+        Content::SubpixelMask => {
+            // Not asked for: `Format::Alpha` is what the renderer requests, so
+            // reaching here means the request changed and this needs writing.
+            // Drawn as a plain mask from the green channel rather than dropped,
+            // so the text is readable rather than absent.
+            for row in 0..h {
+                for col in 0..w {
+                    let i = ((row * w + col) * 4) as usize;
+                    let a = image.data[i + 1];
+                    page.put_pixel(x + col, y + row, image::Rgba([255, 255, 255, a]));
+                }
+            }
+        }
+    }
+}
+
+/// Where a glyph goes on screen, in pixels relative to the text's origin.
+///
+/// `color` is per-glyph (read from the parley run's own `Brush`, since
+/// `UiColor` is parley's brush type here) rather than a single colour for
+/// the whole `RenderedText`, so that a span pushed with its own
+/// `StyleProperty::Brush` (I5's inline rich text: a link, a diff of colour
+/// inside one wrapped paragraph) actually renders in that colour instead of
+/// the buffer's base one.
+#[derive(Clone, Copy)]
+pub struct PlacedGlyph {
+    pub entry: GlyphEntry,
+    pub offset: Vec2,
+    pub color: UiColor,
+}

core/src/render/data.rs

@@ -8,6 +8,15 @@ pub struct WindowUniform {
     pub height: f32,
 }
 
+/// One primitive's placement and what to draw there, in the one arena
+/// every layer shares (`Primitives`). Read from a storage buffer by
+/// **both** shader stages: the vertex stage for the corners of the
+/// primitive it is drawing, the fragment stage for the corners of a
+/// *mask's* primitive, which is generally a different one and often in
+/// another layer. A layer's vertex buffer carries only the slot
+/// ([`instance_slot_layout`]), so there is exactly one copy of a
+/// placement and a mask cannot disagree with what was drawn. See
+/// LAYOUT.md's "Masks with a shape".
 #[repr(C)]
 #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
 pub struct PrimitiveInstance {
@@ -15,25 +24,20 @@ pub struct PrimitiveInstance {
     pub binding: u32,
     pub idx: u32,
     pub mask_idx: MaskIdx,
+    pub move_idx: MoveIdx,
 }
 
-impl PrimitiveInstance {
-    const ATTRIBS: [VertexAttribute; 7] = vertex_attr_array![
-        0 => Float32x2,
-        1 => Float32x2,
-        2 => Float32x2,
-        3 => Float32x2,
-        4 => Uint32,
-        5 => Uint32,
-        6 => Uint32,
-    ];
-
-    pub fn desc() -> VertexBufferLayout<'static> {
-        VertexBufferLayout {
-            array_stride: std::mem::size_of::<Self>() as BufferAddress,
-            step_mode: VertexStepMode::Instance,
-            attributes: &Self::ATTRIBS,
-        }
+/// The vertex layout of a layer's draw order: one `u32` slot into the
+/// global instance arena per instance, stepped per instance. Everything a
+/// primitive is made of used to be here as eight vertex attributes; it
+/// moved into the storage buffer above so the fragment stage can read it
+/// too.
+pub fn instance_slot_layout() -> VertexBufferLayout<'static> {
+    const ATTRIBS: [VertexAttribute; 1] = vertex_attr_array![0 => Uint32];
+    VertexBufferLayout {
+        array_stride: std::mem::size_of::<u32>() as BufferAddress,
+        step_mode: VertexStepMode::Instance,
+        attributes: &ATTRIBS,
     }
 }
 
@@ -43,8 +47,86 @@ impl MaskIdx {
     pub const NONE: Self = Self::preset(u32::MAX);
 }
 
+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
+    /// 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 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
+    /// (`UiRenderState::remove`), so the chain cannot outlive what it
+    /// points at.
+    pub parent: MaskIdx,
+}
+
+/// One widget's cumulative on-screen translation, and the slot of the
+/// ancestor to add on top of it. `parent == u32::MAX` ends the chain. A
+/// pure abs-pixel delta, not a general `UiRegion` remap -- sufficient for
+/// every call site that moves a widget (`ScrollArea`, `Offset`) since both are
+/// translations of an already-drawn subtree. See LAYOUT.md section 2.
+///
+/// `_pad` matches WGSL's storage-buffer layout for `MoveOffset`: `delta` is
+/// a `vec2<f32>`, which gives the struct an 8-byte alignment and rounds its
+/// WGSL size up to 16 bytes even though `delta` + `parent` only total 12 --
+/// the same trap `GlyphPrimitive` documents below. `bytemuck` does not
+/// check this for us, and getting it wrong is a wgpu validation panic at
+/// draw time ("buffer bound ... with size 12 where the shader expects 16"),
+/// not a compile error.
+#[repr(C)]
+#[derive(Debug, Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+pub struct MoveOffset {
+    pub delta: [f32; 2],
+    pub parent: u32,
+    _pad: u32,
+}
+
+impl MoveOffset {
+    pub const NONE_PARENT: u32 = u32::MAX;
+
+    pub fn new(delta: [f32; 2], parent: u32) -> Self {
+        Self {
+            delta,
+            parent,
+            _pad: 0,
+        }
+    }
 }

core/src/render/frame_report.rs

@@ -0,0 +1,557 @@
+use std::time::{Duration, Instant};
+
+/// The frame budget `dumpsys gfxinfo` also uses to call a frame "janky": the
+/// 60Hz vsync period. Kept as the same threshold so a percentage from this
+/// report and a percentage from `gfxinfo` mean the same thing. Only a
+/// fallback now that a caller can read the display's real refresh rate
+/// (`report_at_hz`/`mark_phase`'s callers) -- most devices are 60Hz, but a
+/// 90Hz or 120Hz phone judged against this constant would call every frame
+/// "late" that merely met its own, faster budget.
+pub const JANK_THRESHOLD: Duration = Duration::from_nanos(16_666_667);
+
+/// Enough frames for several minutes of scrolling before the oldest ones
+/// start being overwritten -- the same "diagnostic, not a log" sizing
+/// `FrameStats.kt`'s `CAP` uses on the Compose side, chosen independently
+/// here since a `Duration` is smaller than the six `Long` arrays it keeps.
+/// Bumped from 4096 for RUST.md's "Benchmark v2": a fling+stream+type+
+/// keyboard run is ~6,500+ frames on the Compose side, comfortably under
+/// this so `phase_stats` never has to report a phase as partially evicted.
+const RING_CAPACITY: usize = 16384;
+
+/// One `mark_phase` call: the wall-clock instant and the (0-based,
+/// never-reset-by-`reset`-except-at-`reset`-time) absolute frame index at
+/// which a phase began -- `phase_stats` slices `index_ring` against this to
+/// find which recorded samples belong to which phase, since the ring
+/// itself only keeps the most recent `RING_CAPACITY` samples' *values*,
+/// not which phase they were in.
+struct PhaseMark {
+    name: String,
+    start_index: u64,
+    start_at: Instant,
+}
+
+/// One phase's own slice of a report -- RUST.md's "Benchmark v2" spec's
+/// "per-phase blocks in `FrameReport`... frames, late count/percent...
+/// p50/p90/p99, worst, duration". `Display` matches the shape
+/// `docs/bench/compose-phone-v2-2026-09-06.md`'s report already uses, so
+/// the two apps' reports read the same way side by side.
+pub struct PhaseStats {
+    pub name: String,
+    /// How many frames were recorded during this phase in total -- may
+    /// exceed `late + (samples counted)` if some of this phase's frames
+    /// have since been evicted from the ring by a very long run; that
+    /// case is named in the `Display` rather than silently under-counted.
+    pub frames: u64,
+    pub duration: Duration,
+    pub late: u64,
+    pub late_percent: f64,
+    pub p50: Duration,
+    pub p90: Duration,
+    pub p99: Duration,
+    pub worst: Duration,
+    /// `false` if this phase's frame count exceeds how many samples of it
+    /// are still in the ring -- the percentiles above are then computed
+    /// over whatever survived, not the whole phase. UI_RULES.md: this is
+    /// the "we don't fully know" state, named rather than folded silently
+    /// into a number that looks exact.
+    pub complete: bool,
+}
+
+impl std::fmt::Display for PhaseStats {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        writeln!(
+            f,
+            "  {}: {} frames over {:.1}s{}",
+            self.name,
+            self.frames,
+            self.duration.as_secs_f64(),
+            if self.complete {
+                ""
+            } else {
+                " (ring evicted some of this phase)"
+            },
+        )?;
+        writeln!(f, "    late: {} ({:.1}%)", self.late, self.late_percent)?;
+        writeln!(
+            f,
+            "      total  p50 {:.1}ms  p90 {:.1}ms  p99 {:.1}ms",
+            self.p50.as_secs_f64() * 1000.0,
+            self.p90.as_secs_f64() * 1000.0,
+            self.p99.as_secs_f64() * 1000.0,
+        )?;
+        write!(f, "    worst {:.1}ms", self.worst.as_secs_f64() * 1000.0)
+    }
+}
+
+/// A per-frame wall-time report iris keeps of itself, because `dumpsys
+/// gfxinfo` cannot see a `SurfaceView`'s own GPU-drawn frames at all
+/// (RUST.md's I5 box, "Measurements taken" (b)): it instruments Android's
+/// ordinary Skia/HWUI View-drawing pipeline, which a `wgpu`-rendered
+/// `SurfaceView` bypasses entirely. `record` is meant to be called once per
+/// frame, wrapping the same span Compose's own render report and `gfxinfo`
+/// count -- from the frame's redraw/update start to after the frame is
+/// handed to the platform to present.
+///
+/// **What this does not measure**: wgpu's `present()` call queues the frame
+/// with the compositor and returns; it is not fenced against the GPU
+/// actually finishing the frame or the compositor actually showing it, the
+/// way `gfxinfo`'s own `GPU_DURATION`/vsync accounting is. So a sample here
+/// is "how long the CPU took to build and submit this frame", not
+/// "how long the frame took to reach the screen" -- named in
+/// [`FrameStats`]'s own `Display` line rather than presented as the latter,
+/// per the standing rule against showing an inferred number as a measured
+/// one where the two differ.
+///
+/// Fixed-size ring, no allocation on the hot path -- `report()` is the only
+/// place that allocates (a sort over the current ring), and it is only
+/// ever called from a button tap, not once per frame.
+pub struct FrameReport {
+    ring: Box<[Duration; RING_CAPACITY]>,
+    /// The `submit_to_present` half of each sample in `ring`, same index,
+    /// same lifetime -- kept as a second ring rather than a ring of pairs so
+    /// the existing `ring`/percentile code above is untouched (RUST.md's I5
+    /// "Where iris's frame time goes" CPU/GPU split, added 2026-09-05).
+    /// `ring[i] - submit_ring[i]` is that frame's `redraw_to_submit` half.
+    submit_ring: Box<[Duration; RING_CAPACITY]>,
+    /// The absolute (0-based, since the last `reset`) frame index each
+    /// `ring`/`submit_ring` slot's sample belongs to -- what `phase_stats`
+    /// slices against `PhaseMark::start_index` to tell which recorded
+    /// frames fall in which phase.
+    index_ring: Box<[u64; RING_CAPACITY]>,
+    /// How many of `ring`'s slots hold a real sample -- saturates at
+    /// `RING_CAPACITY`, unlike `total_frames` below which keeps counting.
+    len: usize,
+    pos: usize,
+    /// All frames recorded since the last `reset`, even past `RING_CAPACITY`
+    /// -- what `janky_percent` divides by, so a long run's percentage stays
+    /// correct even once the ring itself only holds the most recent frames.
+    total_frames: u64,
+    janky_frames: u64,
+    /// `mark_phase` calls since the last `reset`, oldest first -- see
+    /// `phase_stats`. Empty on an ordinary run that never calls
+    /// `mark_phase`, so `phase_stats` returns an empty `Vec` and a caller
+    /// prints no "per phase:" section at all, matching RUST.md's "empty/
+    /// absent on an ordinary 'Copy' press, which never marks a phase."
+    phases: Vec<PhaseMark>,
+}
+
+/// One resolved reading. `Display` is the log line both the "Frame report"
+/// button and `transcript-bench.sh`-style scripts read, grep-able on
+/// `"iris frame report"`.
+pub struct FrameStats {
+    pub total_frames: u64,
+    pub janky_percent: f64,
+    pub p50: Duration,
+    pub p90: Duration,
+    pub p99: Duration,
+    pub worst: Duration,
+    /// Median of `redraw_to_submit` -- iris's own CPU work (layout, text,
+    /// primitive building) up to and including building the `queue.submit`
+    /// call, per frame. RUST.md's I5 "Where iris's frame time goes" split,
+    /// added 2026-09-05 to answer "CPU or GPU?" with a number rather than a
+    /// guess.
+    pub cpu_p50: Duration,
+    /// Median of `submit_to_present` -- the `queue.submit` call itself plus
+    /// `present()`, i.e. wherever the driver/GPU/compositor wait actually
+    /// happens. Same caveat as the type's own doc: `present()` is not
+    /// fenced against the GPU actually finishing, so this is "how long the
+    /// CPU was blocked handing the frame off", not the frame's true GPU
+    /// time -- still enough to separate "iris is slow building the frame"
+    /// from "iris is slow handing it to the driver".
+    pub gpu_wait_p50: Duration,
+}
+
+impl std::fmt::Display for FrameStats {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        write!(
+            f,
+            "frames={} janky%={:.2} p50={:.1}ms p90={:.1}ms p99={:.1}ms worst={:.1}ms \
+             (measures redraw-start to after present() is called, not GPU/compositor \
+             completion)",
+            self.total_frames,
+            self.janky_percent,
+            self.p50.as_secs_f64() * 1000.0,
+            self.p90.as_secs_f64() * 1000.0,
+            self.p99.as_secs_f64() * 1000.0,
+            self.worst.as_secs_f64() * 1000.0,
+        )?;
+        write!(
+            f,
+            " cpu_p50={:.1}ms gpu_wait_p50={:.1}ms (redraw-start-to-submit vs. \
+             submit-to-after-present)",
+            self.cpu_p50.as_secs_f64() * 1000.0,
+            self.gpu_wait_p50.as_secs_f64() * 1000.0,
+        )
+    }
+}
+
+impl FrameReport {
+    pub fn new() -> Self {
+        Self {
+            ring: Box::new([Duration::ZERO; RING_CAPACITY]),
+            submit_ring: Box::new([Duration::ZERO; RING_CAPACITY]),
+            index_ring: Box::new([0; RING_CAPACITY]),
+            len: 0,
+            pos: 0,
+            total_frames: 0,
+            janky_frames: 0,
+            phases: Vec::new(),
+        }
+    }
+
+    /// Record one frame's elapsed wall time, with no CPU/GPU split (the
+    /// `submit_to_present` half is recorded as zero, so `cpu_p50` reads as
+    /// the whole frame and `gpu_wait_p50` as nothing -- honest for a caller
+    /// that never measured the split, rather than fabricating one). O(1),
+    /// no allocation.
+    pub fn record(&mut self, elapsed: Duration) {
+        self.record_split(elapsed, Duration::ZERO);
+    }
+
+    /// Record one frame's elapsed wall time, split at `queue.submit`:
+    /// `submit_to_present` is the `queue.submit()` call plus `present()`;
+    /// `total - submit_to_present` is everything before it (layout, text,
+    /// primitive building). RUST.md's I5 "Where iris's frame time goes"
+    /// CPU/GPU split, added 2026-09-05. O(1), no allocation.
+    pub fn record_split(&mut self, total: Duration, submit_to_present: Duration) {
+        self.ring[self.pos] = total;
+        self.submit_ring[self.pos] = submit_to_present;
+        self.index_ring[self.pos] = self.total_frames;
+        self.pos = (self.pos + 1) % RING_CAPACITY;
+        self.len = (self.len + 1).min(RING_CAPACITY);
+        self.total_frames += 1;
+        if total > JANK_THRESHOLD {
+            self.janky_frames += 1;
+        }
+    }
+
+    /// Clears every counter and every sample -- what the "Reset frame
+    /// report" control calls, so a report covers only what was scrolled
+    /// after the button was pressed (the same reason `FrameStats.kt`'s
+    /// `reset()` exists on the Compose side). Also clears every phase
+    /// mark, so a fresh run starts with no "per phase:" section until it
+    /// marks one of its own.
+    pub fn reset(&mut self) {
+        self.len = 0;
+        self.pos = 0;
+        self.total_frames = 0;
+        self.janky_frames = 0;
+        self.phases.clear();
+    }
+
+    /// Marks the start of a named phase at the current moment -- every
+    /// frame recorded from here until the next `mark_phase` (or `reset`)
+    /// belongs to it. RUST.md's "Benchmark v2": a scripted bench run calls
+    /// this once per phase (fling/stream/type/keyboard) so `phase_stats`
+    /// can slice one whole run's frames by what was happening during each.
+    pub fn mark_phase(&mut self, name: &str) {
+        // `phase_stats`'s slicing (`idx >= phase.start_index && idx <
+        // end_index`) silently produces an empty or nonsensical slice for
+        // a phase pushed out of order rather than surfacing the misuse
+        // (docs/REVIEW-2026-09-06.md finding 5).
+        debug_assert!(
+            self.phases
+                .last()
+                .is_none_or(|p| self.total_frames >= p.start_index)
+        );
+        self.phases.push(PhaseMark {
+            name: name.to_string(),
+            start_index: self.total_frames,
+            start_at: Instant::now(),
+        });
+    }
+
+    /// One [`PhaseStats`] per `mark_phase` call since the last `reset`,
+    /// oldest first. `now` closes the last phase's wall-clock span (there
+    /// is no "next phase" instant to use for it); `refresh_hz` is what
+    /// each phase's own `late`/`late_percent` is judged against, read from
+    /// the display rather than assumed -- RUST.md's "Benchmark v2": "late
+    /// count/% against the display's refresh rate."
+    pub fn phase_stats(&self, now: Instant, refresh_hz: f32) -> Vec<PhaseStats> {
+        if self.phases.is_empty() || refresh_hz <= 0.0 {
+            return Vec::new();
+        }
+        let budget = Duration::from_secs_f64(1.0 / refresh_hz as f64);
+        self.phases
+            .iter()
+            .enumerate()
+            .map(|(i, phase)| {
+                let (end_index, end_at) = match self.phases.get(i + 1) {
+                    Some(next) => (next.start_index, next.start_at),
+                    None => (self.total_frames, now),
+                };
+                let frames = end_index.saturating_sub(phase.start_index);
+                let mut samples: Vec<Duration> = (0..self.len)
+                    .filter(|&j| {
+                        let idx = self.index_ring[j];
+                        idx >= phase.start_index && idx < end_index
+                    })
+                    .map(|j| self.ring[j])
+                    .collect();
+                let complete = samples.len() as u64 >= frames;
+                if samples.is_empty() {
+                    return PhaseStats {
+                        name: phase.name.clone(),
+                        frames,
+                        duration: end_at.saturating_duration_since(phase.start_at),
+                        late: 0,
+                        late_percent: 0.0,
+                        p50: Duration::ZERO,
+                        p90: Duration::ZERO,
+                        p99: Duration::ZERO,
+                        worst: Duration::ZERO,
+                        complete,
+                    };
+                }
+                samples.sort_unstable();
+                let pct = |p: usize| samples[(samples.len() * p / 100).min(samples.len() - 1)];
+                let late = samples.iter().filter(|&&d| d > budget).count() as u64;
+                PhaseStats {
+                    name: phase.name.clone(),
+                    frames,
+                    duration: end_at.saturating_duration_since(phase.start_at),
+                    late,
+                    late_percent: 100.0 * late as f64 / samples.len() as f64,
+                    p50: pct(50),
+                    p90: pct(90),
+                    p99: pct(99),
+                    worst: *samples.last().expect("checked not empty above"),
+                    complete,
+                }
+            })
+            .collect()
+    }
+
+    /// `None` if nothing has been recorded since the last reset -- the
+    /// "no frames recorded, scroll first" case, not a zeroed report that
+    /// would read as a real (perfect) measurement.
+    pub fn report(&self) -> Option<FrameStats> {
+        if self.len == 0 {
+            return None;
+        }
+        let mut samples: Vec<Duration> = self.ring[..self.len].to_vec();
+        samples.sort_unstable();
+        let pct = |p: usize| samples[(samples.len() * p / 100).min(samples.len() - 1)];
+
+        // Separate arrays rather than subtracting the two medians above:
+        // medians do not distribute over subtraction, and each needs its
+        // own sort.
+        let submit_samples: Vec<Duration> = self.submit_ring[..self.len].to_vec();
+        let cpu_samples: Vec<Duration> = self.ring[..self.len]
+            .iter()
+            .zip(self.submit_ring[..self.len].iter())
+            .map(|(&total, &submit_to_present)| total.saturating_sub(submit_to_present))
+            .collect();
+        let median = |mut v: Vec<Duration>| {
+            v.sort_unstable();
+            v[v.len() / 2]
+        };
+
+        Some(FrameStats {
+            total_frames: self.total_frames,
+            janky_percent: 100.0 * self.janky_frames as f64 / self.total_frames as f64,
+            p50: pct(50),
+            p90: pct(90),
+            p99: pct(99),
+            worst: *samples.last().expect("len > 0 checked above"),
+            cpu_p50: median(cpu_samples),
+            gpu_wait_p50: median(submit_samples),
+        })
+    }
+
+    /// `(late count, late percent)` over every sample still in the ring,
+    /// judged against `refresh_hz`'s own frame budget rather than the
+    /// fixed 60Hz `JANK_THRESHOLD` -- RUST.md's "Benchmark v2": "late
+    /// count/% against the display's refresh rate... print 'at N Hz (X ms
+    /// budget)' like Compose does." A separate method from `report()`
+    /// rather than a parameter on it, so `report()`'s own `janky_percent`
+    /// (and the exact-boundary test pinned to `JANK_THRESHOLD`) is
+    /// unaffected for every existing caller that never measured a real
+    /// refresh rate. `(0, 0.0)` with nothing recorded or a non-positive
+    /// `refresh_hz`.
+    pub fn late_at_hz(&self, refresh_hz: f32) -> (u64, f64) {
+        if self.len == 0 || refresh_hz <= 0.0 {
+            return (0, 0.0);
+        }
+        let budget = Duration::from_secs_f64(1.0 / refresh_hz as f64);
+        let late = self.ring[..self.len]
+            .iter()
+            .filter(|&&d| d > budget)
+            .count() as u64;
+        (late, 100.0 * late as f64 / self.len as f64)
+    }
+}
+
+impl Default for FrameReport {
+    fn default() -> Self {
+        Self::new()
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn no_frames_reports_none() {
+        assert!(FrameReport::new().report().is_none());
+    }
+
+    #[test]
+    fn one_frame_is_every_percentile_and_the_worst() {
+        let mut r = FrameReport::new();
+        r.record(Duration::from_millis(10));
+        let stats = r.report().unwrap();
+        assert_eq!(stats.total_frames, 1);
+        assert_eq!(stats.p50, Duration::from_millis(10));
+        assert_eq!(stats.p99, Duration::from_millis(10));
+        assert_eq!(stats.worst, Duration::from_millis(10));
+        assert_eq!(stats.janky_percent, 0.0);
+    }
+
+    #[test]
+    fn percentiles_and_worst_over_a_known_set() {
+        let mut r = FrameReport::new();
+        // 100 samples, 1ms..=100ms, fed out of order so the ring's own
+        // order is not what gives the right answer -- the sort has to.
+        for ms in (1..=100).rev() {
+            r.record(Duration::from_millis(ms));
+        }
+        let stats = r.report().unwrap();
+        assert_eq!(stats.total_frames, 100);
+        assert_eq!(stats.p50, Duration::from_millis(51));
+        assert_eq!(stats.p90, Duration::from_millis(91));
+        assert_eq!(stats.p99, Duration::from_millis(100));
+        assert_eq!(stats.worst, Duration::from_millis(100));
+    }
+
+    #[test]
+    fn jank_threshold_matches_gfxinfos_60hz_budget() {
+        let mut r = FrameReport::new();
+        r.record(Duration::from_nanos(16_666_667)); // exactly on budget: not janky
+        r.record(Duration::from_nanos(16_666_668)); // one ns over: janky
+        let stats = r.report().unwrap();
+        assert_eq!(stats.janky_percent, 50.0);
+    }
+
+    #[test]
+    fn janky_percent_is_over_all_time_frames_not_just_the_ring() {
+        // Fewer than RING_CAPACITY frames, all janky, then a fresh reset --
+        // the percentage must reset to 0, not divide by a stale count.
+        let mut r = FrameReport::new();
+        for _ in 0..10 {
+            r.record(Duration::from_millis(50));
+        }
+        assert_eq!(r.report().unwrap().janky_percent, 100.0);
+        r.reset();
+        assert!(r.report().is_none());
+        r.record(Duration::from_millis(1));
+        assert_eq!(r.report().unwrap().janky_percent, 0.0);
+    }
+
+    #[test]
+    fn record_without_a_split_reports_the_whole_frame_as_cpu() {
+        // A caller that never measured the split (plain `record`) should
+        // not fabricate a GPU-wait number -- it reads as zero, and the CPU
+        // half reads as the whole frame.
+        let mut r = FrameReport::new();
+        r.record(Duration::from_millis(20));
+        let stats = r.report().unwrap();
+        assert_eq!(stats.cpu_p50, Duration::from_millis(20));
+        assert_eq!(stats.gpu_wait_p50, Duration::ZERO);
+    }
+
+    #[test]
+    fn record_split_reports_each_halfs_own_median() {
+        let mut r = FrameReport::new();
+        // Three frames: total is always 30ms, but the CPU/GPU-wait split
+        // moves, so the two medians must be independent of each other and
+        // of `total`'s own median.
+        r.record_split(Duration::from_millis(30), Duration::from_millis(5));
+        r.record_split(Duration::from_millis(30), Duration::from_millis(10));
+        r.record_split(Duration::from_millis(30), Duration::from_millis(20));
+        let stats = r.report().unwrap();
+        assert_eq!(stats.p50, Duration::from_millis(30));
+        assert_eq!(stats.gpu_wait_p50, Duration::from_millis(10));
+        assert_eq!(stats.cpu_p50, Duration::from_millis(20));
+    }
+
+    #[test]
+    fn ring_wraps_without_growing_past_capacity() {
+        let mut r = FrameReport::new();
+        for i in 0..(RING_CAPACITY * 2) {
+            r.record(Duration::from_millis(1 + (i % 5) as u64));
+        }
+        let stats = r.report().unwrap();
+        // total_frames keeps the full count even once the ring has wrapped.
+        assert_eq!(stats.total_frames, (RING_CAPACITY * 2) as u64);
+        // but every sample the ring can report on is still one of the five
+        // values fed in, since a wrap can only overwrite with more of the
+        // same pattern here.
+        assert!(stats.worst <= Duration::from_millis(5));
+    }
+
+    #[test]
+    fn no_marks_means_no_phases() {
+        let mut r = FrameReport::new();
+        r.record(Duration::from_millis(5));
+        assert!(r.phase_stats(Instant::now(), 60.0).is_empty());
+    }
+
+    #[test]
+    fn phases_slice_frames_by_when_they_were_marked() {
+        let mut r = FrameReport::new();
+        r.mark_phase("a");
+        for _ in 0..5 {
+            r.record(Duration::from_millis(10)); // 10ms: late at 60Hz (16.7ms budget)... no, 10<16.7, not late
+        }
+        r.mark_phase("b");
+        for _ in 0..3 {
+            r.record(Duration::from_millis(20)); // 20ms: late at 60Hz
+        }
+        let now = Instant::now();
+        let phases = r.phase_stats(now, 60.0);
+        assert_eq!(phases.len(), 2);
+        assert_eq!(phases[0].name, "a");
+        assert_eq!(phases[0].frames, 5);
+        assert_eq!(phases[0].late, 0);
+        assert_eq!(phases[0].worst, Duration::from_millis(10));
+        assert_eq!(phases[1].name, "b");
+        assert_eq!(phases[1].frames, 3);
+        assert_eq!(phases[1].late, 3);
+        assert_eq!(phases[1].late_percent, 100.0);
+        assert_eq!(phases[1].worst, Duration::from_millis(20));
+        assert!(phases[0].complete);
+        assert!(phases[1].complete);
+    }
+
+    #[test]
+    fn the_last_phase_runs_until_now() {
+        let mut r = FrameReport::new();
+        r.mark_phase("only");
+        r.record(Duration::from_millis(1));
+        std::thread::sleep(Duration::from_millis(20));
+        let now = Instant::now();
+        let phases = r.phase_stats(now, 60.0);
+        assert_eq!(phases.len(), 1);
+        assert!(phases[0].duration >= Duration::from_millis(20));
+    }
+
+    #[test]
+    fn reset_clears_phase_marks() {
+        let mut r = FrameReport::new();
+        r.mark_phase("a");
+        r.record(Duration::from_millis(1));
+        r.reset();
+        assert!(r.phase_stats(Instant::now(), 60.0).is_empty());
+    }
+
+    #[test]
+    fn late_at_hz_uses_the_given_refresh_rate_not_the_fixed_60hz_constant() {
+        let mut r = FrameReport::new();
+        // 10ms is under 60Hz's 16.7ms budget but over 120Hz's 8.3ms one.
+        r.record(Duration::from_millis(10));
+        assert_eq!(r.late_at_hz(60.0), (0, 0.0));
+        assert_eq!(r.late_at_hz(120.0), (1, 100.0));
+    }
+}

core/src/render/mod.rs

@@ -1,30 +1,141 @@
-use std::num::NonZero;
-
 use crate::{
     UiData, UiRenderState,
-    render::{data::PrimitiveInstance, texture::GpuTextures, util::ArrBuf},
-    util::HashMap,
+    render::{
+        data::{PrimitiveInstance, instance_slot_layout},
+        texture::GpuTextures,
+        util::ArrBuf,
+    },
+    util::{HashMap, Vec2},
 };
 use data::WindowUniform;
+use pollster::FutureExt;
 use wgpu::{
     util::{BufferInitDescriptor, DeviceExt},
     *,
 };
-use winit::dpi::PhysicalSize;
 
+mod atlas;
 mod data;
+mod frame_report;
 mod primitive;
+mod sdf;
 mod texture;
 mod util;
 
-pub use data::{Mask, MaskIdx};
+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
+/// `Adapter::request_device` for -- shared so the two copies cannot drift,
+/// per AGENTS.md's "write the logic once."
+///
+/// Built from `Limits::default()`, **not** a downlevel variant: the shader
+/// (`shader.wgsl`) reads four `var<storage>` buffers (rects, glyphs, masks,
+/// move_offsets) from the vertex stage, and `Limits::downlevel_webgl2_defaults()`
+/// zeroes `max_storage_buffers_per_shader_stage` along with the compute
+/// limits below -- switching to it would trade one `request_device` crash
+/// for a bind-group-layout one on the same downlevel hardware this is meant
+/// to support. `max_buffer_size` is raised for the growing instance/atlas
+/// buffers (`ArrBuf`, `GpuTextures`); everything else is `default()`'s
+/// desktop-tier value, unchanged.
+///
+/// The six `max_compute_*` fields are zeroed because nothing in this crate
+/// creates a `ComputePipeline` or writes a `@compute` shader stage --
+/// grepped for both across `iris`/`iris-core` before writing this, found
+/// none. `Limits::default()` requests desktop-tier compute limits
+/// unconditionally (`max_compute_workgroups_per_dimension: 65535`) even
+/// though nothing asks a device to actually support compute, which is what
+/// crashed `request_device` on the Android emulator's software GL path
+/// (`EMU_GPU=software`, `force-gles`): SwiftShader's GL reports itself as
+/// OpenGL ES 3.0, which has no compute shaders, so the adapter's real limit
+/// is 0 and the unconditional request fails outright
+/// (`RUST.md`'s "Software mode ... crashes for a third, different reason").
+/// The same would happen on a real GLES-3.0-only Android device. If a
+/// future change adds a compute pass, request the specific limits it needs
+/// here rather than reverting to the desktop-tier default for everything.
+pub fn device_limits() -> Limits {
+    Limits {
+        max_buffer_size: 1 << 30,
+        max_compute_workgroup_storage_size: 0,
+        max_compute_invocations_per_workgroup: 0,
+        max_compute_workgroup_size_x: 0,
+        max_compute_workgroup_size_y: 0,
+        max_compute_workgroup_size_z: 0,
+        max_compute_workgroups_per_dimension: 0,
+        ..Default::default()
+    }
+}
+
+/// A capped log of wgpu's *uncaptured* errors -- everything that reaches
+/// `Device::on_uncaptured_error` rather than one of `UiRenderNode::new`'s
+/// own error scopes, i.e. every wgpu error raised outside device/pipeline
+/// creation: a validation failure during an ordinary frame's `update`/
+/// `draw`, for instance. wgpu's default handler for these is `panic!` with
+/// no caller able to intervene -- exactly what aborted the P0 bench APK
+/// once already (this file's `UiRenderNode::new` doc comment) -- so both
+/// platform backends install a handler here instead of leaving the default
+/// in place, per RUST.md's P0 box ("every wgpu uncaptured error ... it
+/// must never panic in release").
+///
+/// Cheap to `Clone` (an `Arc` around the real storage) rather than a
+/// process-wide static, so a caller builds one alongside its `Device`,
+/// hands one clone to `on_uncaptured_error`'s closure and keeps the other
+/// for the Diagnostics page to read -- context passed explicitly, per
+/// AGENTS.md/CODE_RULES.md's "no globals" rather than reached for through a
+/// `OnceLock`.
+#[derive(Clone)]
+pub struct WgpuErrorLog {
+    errors: std::sync::Arc<std::sync::Mutex<std::collections::VecDeque<String>>>,
+}
+
+/// How many uncaptured errors the log keeps -- old ones drop off the front
+/// rather than being trimmed on read, so a build spraying errors every
+/// frame doesn't grow this without bound.
+const WGPU_ERROR_LOG_CAP: usize = 20;
+
+impl Default for WgpuErrorLog {
+    fn default() -> Self {
+        Self {
+            errors: std::sync::Arc::new(std::sync::Mutex::new(std::collections::VecDeque::new())),
+        }
+    }
+}
+
+impl WgpuErrorLog {
+    pub fn record(&self, error: impl std::fmt::Display) {
+        let mut errors = self.errors.lock().unwrap();
+        if errors.len() >= WGPU_ERROR_LOG_CAP {
+            errors.pop_front();
+        }
+        errors.push_back(error.to_string());
+    }
+
+    /// A snapshot for the Diagnostics page -- cloned rather than held,
+    /// since the lock must not outlive one call.
+    pub fn snapshot(&self) -> Vec<String> {
+        self.errors.lock().unwrap().iter().cloned().collect()
+    }
+}
 
 pub struct UiRenderNode {
     uniform_group: BindGroup,
     primitive_layout: BindGroupLayout,
+    /// Group 1: `rects` and `glyphs`. Global and bound once per frame,
+    /// not per layer -- a mask referencing a rect drawn in another layer
+    /// has to be able to read it (see `Primitives`).
+    primitives: PrimitiveBuffers,
+    primitive_group: BindGroup,
     rsc_layout: BindGroupLayout,
     rsc_group: BindGroup,
 
@@ -34,28 +145,76 @@ pub struct UiRenderNode {
     active: Vec<usize>,
     window_buffer: Buffer,
     textures: GpuTextures,
+    /// Every primitive's placement, read by the vertex stage for the
+    /// primitive being drawn and by the fragment stage for a mask's.
+    instances: ArrBuf<PrimitiveInstance>,
     masks: ArrBuf<Mask>,
+    move_offsets: ArrBuf<MoveOffset>,
+    /// Group 3: the masks and move-offsets storage buffers, on their own --
+    /// see IRIS_TODO.md's "Appending one image ... rebuilds every other
+    /// image's bind group". These used to live in group 2 alongside each
+    /// standalone image's own texture view, so an image's bind group named
+    /// the masks/move_offsets buffer directly; the moment either buffer
+    /// resized (which a widget getting its *first* move slot can trigger,
+    /// unrelated to any image), `ArrBuf::update` handed back a new `Buffer`
+    /// identity and every image's bind group -- one per live image -- had
+    /// to be rebuilt to reference it. Pulling both buffers into their own
+    /// group, bound once per frame rather than once per draw call, means a
+    /// buffer resize now rebuilds exactly this one group instead of N.
+    masks_layout: BindGroupLayout,
+    masks_group: BindGroup,
 }
 
+/// One layer's vertex buffers: the slots it draws, in order. The
+/// primitives themselves are in `UiRenderNode::instances`.
 struct RenderLayer {
-    instance: ArrBuf<PrimitiveInstance>,
-    primitives: PrimitiveBuffers,
-    primitive_group: BindGroup,
+    order: ArrBuf<u32>,
+    /// A standalone image's slots, kept apart from `order` because each
+    /// one draws with its own bind group -- see `UiRenderNode::draw`.
+    images: ArrBuf<u32>,
+    /// The texture slot each entry of `images` draws with, in the same
+    /// order, refreshed alongside it. Not in the vertex buffer itself
+    /// because it names a bind group, not shader data.
+    image_tex_indices: Vec<u32>,
 }
 
 impl UiRenderNode {
     pub fn draw<'a>(&'a self, pass: &mut RenderPass<'a>) {
         pass.set_pipeline(&self.pipeline);
         pass.set_bind_group(0, &self.uniform_group, &[]);
-        pass.set_bind_group(2, &self.rsc_group, &[]);
+        // Group 1 is global now, so it is set here rather than per layer.
+        pass.set_bind_group(1, &self.primitive_group, &[]);
+        // Set once, not per layer or per image: masks/move_offsets are read
+        // by every primitive and every standalone image alike, and living
+        // in their own group (rather than folded into group 2 alongside the
+        // per-image texture view) is what keeps an image's own bind group
+        // from naming a buffer that changes size on an unrelated widget's
+        // first draw -- see the comment on `masks_group` below.
+        pass.set_bind_group(3, &self.masks_group, &[]);
         for i in &self.active {
             let layer = &self.layers[i];
-            if layer.instance.len() == 0 {
+            if layer.order.len() == 0 && layer.images.len() == 0 {
                 continue;
             }
-            pass.set_bind_group(1, &layer.primitive_group, &[]);
-            pass.set_vertex_buffer(0, layer.instance.buffer.slice(..));
-            pass.draw(0..4, 0..layer.instance.len() as u32);
+            if layer.order.len() > 0 {
+                pass.set_bind_group(2, &self.rsc_group, &[]);
+                pass.set_vertex_buffer(0, layer.order.buffer.slice(..));
+                pass.draw(0..4, 0..layer.order.len() as u32);
+            }
+            // Images draw after this layer's rects and glyphs, one draw call
+            // each with its own bind group. That draws every image "on top"
+            // within the layer, which loses nothing that currently exists:
+            // `Primitives::apply_free` frees with `swap_remove`, so a layer's
+            // draw order was already undefined before images had their own
+            // list -- nothing before this relied on interleaving a rect
+            // between two images at a particular position.
+            if layer.images.len() > 0 {
+                pass.set_vertex_buffer(0, layer.images.buffer.slice(..));
+                for (k, &tex_idx) in layer.image_tex_indices.iter().enumerate() {
+                    pass.set_bind_group(2, self.textures.image_bind_group(tex_idx), &[]);
+                    pass.draw(0..4, k as u32..k as u32 + 1);
+                }
+            }
         }
     }
 
@@ -65,79 +224,156 @@ impl UiRenderNode {
         queue: &Queue,
         ui: &mut UiData,
         ui_render: &mut UiRenderState,
-    ) {
+    ) -> FrameUpdateStats {
         self.active.clear();
-        for (i, primitives) in ui_render.layers.iter_mut() {
+        for (i, order) in ui_render.layers.iter_mut() {
             self.active.push(i);
-            for change in primitives.apply_free() {
-                if let Some(inst) = ui_render.active.get_mut(&change.id) {
-                    for h in &mut inst.primitives {
-                        if h.layer == i && h.inst_idx == change.old {
-                            h.inst_idx = change.new;
-                            break;
-                        }
-                    }
-                }
-            }
-            let rlayer = self.layers.entry(i).or_insert_with(|| {
-                let primitives = PrimitiveBuffers::new(device);
-                let primitive_group =
-                    Self::primitive_group(device, &self.primitive_layout, primitives.buffers());
-                RenderLayer {
-                    instance: ArrBuf::new(
-                        device,
-                        BufferUsages::VERTEX | BufferUsages::COPY_DST,
-                        "instance",
-                    ),
-                    primitives,
-                    primitive_group,
-                }
-            });
-            if primitives.updated {
-                rlayer
-                    .instance
-                    .update(device, queue, primitives.instances());
-                rlayer.primitives.update(device, queue, primitives.data());
-                rlayer.primitive_group = Self::primitive_group(
+            let rlayer = self.layers.entry(i).or_insert_with(|| RenderLayer {
+                order: ArrBuf::new(
                     device,
-                    &self.primitive_layout,
-                    rlayer.primitives.buffers(),
-                );
-                primitives.updated = false;
+                    BufferUsages::VERTEX | BufferUsages::COPY_DST,
+                    "layer order",
+                ),
+                images: ArrBuf::new(
+                    device,
+                    BufferUsages::VERTEX | BufferUsages::COPY_DST,
+                    "layer image order",
+                ),
+                image_tex_indices: Vec::new(),
+            });
+            if order.updated {
+                rlayer.order.update(device, queue, order.order());
+                rlayer.images.update(device, queue, order.images());
+                rlayer.image_tex_indices = order
+                    .images()
+                    .iter()
+                    .map(|&slot| ui_render.primitives.instance(slot).idx)
+                    .collect();
+                order.updated = false;
             }
         }
-        let mut changed = false;
-        changed |= self.textures.update(&mut ui.textures);
-        if ui.masks.changed {
+        let instances_resized = if ui_render.primitives.updated {
+            ui_render.primitives.updated = false;
+            let resized = self
+                .instances
+                .update(device, queue, ui_render.primitives.instances());
+            self.primitives
+                .update(device, queue, ui_render.primitives.data());
+            self.primitive_group =
+                Self::primitive_group(device, &self.primitive_layout, self.primitives.buffers());
+            resized
+        } else {
+            false
+        };
+        let masks_resized = if ui.masks.changed {
             ui.masks.changed = false;
-            self.masks.update(device, queue, &ui.masks[..]);
-            changed = true;
+            self.masks.update(device, queue, &ui.masks[..])
+        } else {
+            false
+        };
+        let moves_resized = if ui.move_offsets.changed {
+            ui.move_offsets.changed = false;
+            self.move_offsets
+                .update(device, queue, &ui.move_offsets[..])
+        } else {
+            false
+        };
+        if masks_resized || moves_resized || instances_resized {
+            self.masks_group = Self::masks_group(
+                device,
+                &self.masks_layout,
+                &self.masks,
+                &self.move_offsets,
+                &self.instances,
+            );
         }
-        if changed {
-            self.rsc_group = Self::rsc_group(device, &self.rsc_layout, &self.textures, &self.masks);
+        let rebuild_main = self.textures.update(&mut ui.textures, &self.rsc_layout);
+        if rebuild_main {
+            self.rsc_group = Self::rsc_group(device, &self.rsc_layout, &self.textures);
+        }
+        FrameUpdateStats {
+            masks_resized,
+            moves_resized,
         }
     }
 
-    pub fn resize(&mut self, size: &PhysicalSize<u32>, queue: &Queue) {
+    /// Takes a size rather than a window type: this is the only thing the
+    /// core wanted from winit, and depending on a windowing backend for two
+    /// numbers is what put `android-activity` in the core's graph for an
+    /// Android build that is meant to go through android-view instead.
+    pub fn resize(&mut self, size: impl Into<Vec2>, queue: &Queue) {
+        let size = size.into();
         let slice = &[WindowUniform {
-            width: size.width as f32,
-            height: size.height as f32,
+            width: size.x,
+            height: size.y,
         }];
         queue.write_buffer(&self.window_buffer, 0, bytemuck::cast_slice(slice));
     }
 
+    /// Builds every bind group layout, the pipeline, and the two storage
+    /// buffers this needs -- fallibly, since this is exactly the call that
+    /// aborted the process on Iris's phone in a release build with no
+    /// message beyond "wgpu error: Validation Error" (RUST.md's P0 box,
+    /// "iris bench crash on the phone, 2026-09-06"). wgpu's own default
+    /// behaviour for an uncaptured error is `panic!` with no caller able to
+    /// intervene, so every `create_bind_group_layout`/`create_render_pipeline`
+    /// call below runs inside three nested error scopes (one per
+    /// `ErrorFilter`) instead: whichever scope catches something, its
+    /// `wgpu::Error`'s `Display` is wgpu-core's own `format_error` output
+    /// (`"Validation Error\n\nCaused by:\n  ..."`, the same text the panic
+    /// would have printed before Android's crash reporter truncated it) and
+    /// becomes this function's `Err`. Both callers
+    /// (`android::render::AndroidRenderer::new`, `default::render::
+    /// UiRenderer::new`) already call `Device`-creation with
+    /// `pollster::block_on`, so returning a plain `Result` here rather than
+    /// making this `async fn` keeps that same synchronous shape.
     pub fn new(
         device: &Device,
         queue: &Queue,
         config: &SurfaceConfiguration,
-        limits: UiLimits,
-    ) -> Self {
+        window_size: impl Into<Vec2>,
+    ) -> Result<Self, String> {
+        // Popped in reverse of this order, once every creation call below
+        // has run -- `Device::push_error_scope`'s own contract.
+        let oom_scope = device.push_error_scope(ErrorFilter::OutOfMemory);
+        let validation_scope = device.push_error_scope(ErrorFilter::Validation);
+        let internal_scope = device.push_error_scope(ErrorFilter::Internal);
+
         let shader = device.create_shader_module(ShaderModuleDescriptor {
             label: Some("UI Shape Shader"),
             source: ShaderSource::Wgsl(SHAPE_SHADER.into()),
         });
 
-        let window_uniform = WindowUniform::default();
+        // Seeded from the caller's own reported size, not
+        // `WindowUniform::default()` (0, 0): the vertex shader divides by
+        // `window.dim` to reach clip space, so a window this buffer
+        // disagrees with means every primitive's position is NaN/Inf and is
+        // dropped before rasterization -- the clear colour still reaches
+        // the screen (the pass runs regardless) while nothing drawn on top
+        // of it ever does. winit's backend gets away with the old default
+        // because winit fires an initial `WindowEvent::Resized` that calls
+        // `resize()` before the first frame; android-view has no such
+        // automatic event, so `AndroidRenderer::new` built a node whose
+        // window buffer was never corrected -- this is I2's "nothing draws"
+        // bug (RUST.md).
+        //
+        // **Deliberately not `config.width`/`config.height`**: those are
+        // the surface's *physical* pixel size, which the swapchain needs,
+        // but everything downstream of this uniform (layout, hit-testing,
+        // glyph/rect positions) works in the caller's own units -- on
+        // Android that's *logical* (physical / density) since RUST.md's P0
+        // box ("text is far too small"), on desktop it's whatever
+        // `default::render::UiRenderer::new` already divides by
+        // `window.scale_factor()`. Passing it in explicitly, rather than
+        // deriving it from `config` here, is what keeps this crate from
+        // needing to know either platform's notion of density at all.
+        let window_uniform = {
+            let size = window_size.into();
+            WindowUniform {
+                width: size.x,
+                height: size.y,
+            }
+        };
         let window_buffer = device.create_buffer_init(&BufferInitDescriptor {
             label: Some("window"),
             contents: bytemuck::cast_slice(&[window_uniform]),
@@ -161,34 +397,53 @@ impl UiRenderNode {
         let uniform_group = Self::bind_group_0(device, &uniform_layout, &window_buffer);
 
         let primitive_layout = device.create_bind_group_layout(&BindGroupLayoutDescriptor {
-            entries: &core::array::from_fn::<_, { PrimitiveBuffers::LEN }, _>(|i| {
-                BindGroupLayoutEntry {
-                    binding: i as u32,
-                    visibility: ShaderStages::FRAGMENT,
-                    ty: BindingType::Buffer {
-                        ty: BufferBindingType::Storage { read_only: true },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                    count: None,
-                }
+            entries: &PrimitiveBuffers::BINDINGS.map(|binding| BindGroupLayoutEntry {
+                binding,
+                visibility: ShaderStages::FRAGMENT,
+                ty: BindingType::Buffer {
+                    ty: BufferBindingType::Storage { read_only: true },
+                    has_dynamic_offset: false,
+                    min_binding_size: None,
+                },
+                count: None,
             }),
             label: Some("primitive"),
         });
 
         let tex_manager = GpuTextures::new(device, queue);
+        let primitives = PrimitiveBuffers::new(device);
+        let primitive_group =
+            Self::primitive_group(device, &primitive_layout, primitives.buffers());
+        let instances = ArrBuf::new(
+            device,
+            BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            "ui instances",
+        );
         let masks = ArrBuf::new(
             device,
             BufferUsages::STORAGE | BufferUsages::COPY_DST,
             "ui masks",
         );
+        let move_offsets = ArrBuf::new(
+            device,
+            BufferUsages::STORAGE | BufferUsages::COPY_DST,
+            "ui move offsets",
+        );
 
-        let rsc_layout = Self::rsc_layout(device, &limits);
-        let rsc_group = Self::rsc_group(device, &rsc_layout, &tex_manager, &masks);
+        let rsc_layout = Self::rsc_layout(device);
+        let rsc_group = Self::rsc_group(device, &rsc_layout, &tex_manager);
+        let masks_layout = Self::masks_layout(device);
+        let masks_group =
+            Self::masks_group(device, &masks_layout, &masks, &move_offsets, &instances);
 
         let pipeline_layout = device.create_pipeline_layout(&PipelineLayoutDescriptor {
             label: Some("UI Shape Pipeline Layout"),
-            bind_group_layouts: &[&uniform_layout, &primitive_layout, &rsc_layout],
+            bind_group_layouts: &[
+                Some(&uniform_layout),
+                Some(&primitive_layout),
+                Some(&rsc_layout),
+                Some(&masks_layout),
+            ],
             immediate_size: 0,
         });
         let pipeline = device.create_render_pipeline(&RenderPipelineDescriptor {
@@ -197,7 +452,7 @@ impl UiRenderNode {
             vertex: VertexState {
                 module: &shader,
                 entry_point: Some("vs_main"),
-                buffers: &[PrimitiveInstance::desc()],
+                buffers: &[Some(instance_slot_layout())],
                 compilation_options: Default::default(),
             },
             fragment: Some(FragmentState {
@@ -229,9 +484,22 @@ impl UiRenderNode {
             cache: None,
         });
 
-        Self {
+        // Reverse of the push order above. Only one of these should ever be
+        // `Some` in practice -- three separate scopes exist to name *which*
+        // kind of error it was, not because more than one is expected at
+        // once.
+        let internal_err = internal_scope.pop().block_on();
+        let validation_err = validation_scope.pop().block_on();
+        let oom_err = oom_scope.pop().block_on();
+        if let Some(err) = validation_err.or(oom_err).or(internal_err) {
+            return Err(err.to_string());
+        }
+
+        Ok(Self {
             uniform_group,
             primitive_layout,
+            primitives,
+            primitive_group,
             rsc_layout,
             rsc_group,
             pipeline,
@@ -239,8 +507,12 @@ impl UiRenderNode {
             layers: HashMap::default(),
             active: Vec::new(),
             textures: tex_manager,
+            instances,
             masks,
-        }
+            move_offsets,
+            masks_layout,
+            masks_group,
+        })
     }
 
     fn bind_group_0(
@@ -273,7 +545,14 @@ impl UiRenderNode {
         })
     }
 
-    fn rsc_layout(device: &Device, limits: &UiLimits) -> BindGroupLayout {
+    /// Group 2: the shared atlas array and one standalone-image slot (a null
+    /// view for the main draw, a real one for each image's own bind group --
+    /// see `GpuTextures`), plus one sampler. No `count` on any entry: this
+    /// needs nothing beyond plain Vulkan 1.0 / GLES sampling, unlike the
+    /// `binding_array` layout it replaced (see TEXTURES.md's "Recommended
+    /// shape"). Masks and move_offsets are deliberately *not* here -- see
+    /// `masks_layout` below for why they get their own group.
+    fn rsc_layout(device: &Device) -> BindGroupLayout {
         device.create_bind_group_layout(&BindGroupLayoutDescriptor {
             entries: &[
                 BindGroupLayoutEntry {
@@ -281,20 +560,91 @@ impl UiRenderNode {
                     visibility: ShaderStages::FRAGMENT,
                     ty: BindingType::Texture {
                         sample_type: TextureSampleType::Float { filterable: false },
-                        view_dimension: TextureViewDimension::D2,
+                        view_dimension: TextureViewDimension::D2Array,
                         multisampled: false,
                     },
-                    count: Some(NonZero::new(limits.max_textures).unwrap()),
+                    count: None,
                 },
                 BindGroupLayoutEntry {
                     binding: 1,
                     visibility: ShaderStages::FRAGMENT,
-                    ty: BindingType::Sampler(SamplerBindingType::NonFiltering),
-                    count: Some(NonZero::new(limits.max_samplers).unwrap()),
+                    ty: BindingType::Texture {
+                        sample_type: TextureSampleType::Float { filterable: false },
+                        view_dimension: TextureViewDimension::D2,
+                        multisampled: false,
+                    },
+                    count: None,
                 },
                 BindGroupLayoutEntry {
                     binding: 2,
                     visibility: ShaderStages::FRAGMENT,
+                    ty: BindingType::Sampler(SamplerBindingType::NonFiltering),
+                    count: None,
+                },
+            ],
+            label: Some("ui rsc"),
+        })
+    }
+
+    /// The main group: rects and glyphs never sample the image slot, so it
+    /// gets a 1x1 null view rather than any live standalone image's.
+    fn rsc_group(
+        device: &Device,
+        layout: &BindGroupLayout,
+        tex_manager: &GpuTextures,
+    ) -> BindGroup {
+        device.create_bind_group(&BindGroupDescriptor {
+            layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: BindingResource::TextureView(tex_manager.array_view()),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: BindingResource::TextureView(tex_manager.null_view()),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: BindingResource::Sampler(tex_manager.sampler()),
+                },
+            ],
+            label: Some("ui rsc"),
+        })
+    }
+
+    /// Group 3: the masks and move_offsets storage buffers, shared by the
+    /// main draw and every standalone image alike (see the field comment on
+    /// `masks_group`). Bound once per frame in `draw()` rather than folded
+    /// into group 2, so a resize of either buffer -- which an unrelated
+    /// widget's first move slot can trigger -- rebuilds this one group
+    /// instead of every image's.
+    fn masks_layout(device: &Device) -> BindGroupLayout {
+        device.create_bind_group_layout(&BindGroupLayoutDescriptor {
+            entries: &[
+                BindGroupLayoutEntry {
+                    binding: 0,
+                    visibility: ShaderStages::FRAGMENT,
+                    ty: BindingType::Buffer {
+                        ty: BufferBindingType::Storage { read_only: true },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                BindGroupLayoutEntry {
+                    binding: 1,
+                    visibility: ShaderStages::VERTEX | ShaderStages::FRAGMENT,
+                    ty: BindingType::Buffer {
+                        ty: BufferBindingType::Storage { read_only: true },
+                        has_dynamic_offset: false,
+                        min_binding_size: None,
+                    },
+                    count: None,
+                },
+                BindGroupLayoutEntry {
+                    binding: 2,
+                    visibility: ShaderStages::VERTEX | ShaderStages::FRAGMENT,
                     ty: BindingType::Buffer {
                         ty: BufferBindingType::Storage { read_only: true },
                         has_dynamic_offset: false,
@@ -303,60 +653,67 @@ impl UiRenderNode {
                     count: None,
                 },
             ],
-            label: Some("ui rsc"),
+            label: Some("ui masks"),
         })
     }
 
-    fn rsc_group(
+    fn masks_group(
         device: &Device,
         layout: &BindGroupLayout,
-        tex_manager: &GpuTextures,
         masks: &ArrBuf<Mask>,
+        move_offsets: &ArrBuf<MoveOffset>,
+        instances: &ArrBuf<PrimitiveInstance>,
     ) -> BindGroup {
         device.create_bind_group(&BindGroupDescriptor {
             layout,
             entries: &[
                 BindGroupEntry {
                     binding: 0,
-                    resource: BindingResource::TextureViewArray(&tex_manager.views()),
+                    resource: masks.buffer.as_entire_binding(),
                 },
                 BindGroupEntry {
                     binding: 1,
-                    resource: BindingResource::SamplerArray(&tex_manager.samplers()),
+                    resource: move_offsets.buffer.as_entire_binding(),
                 },
                 BindGroupEntry {
                     binding: 2,
-                    resource: masks.buffer.as_entire_binding(),
+                    resource: instances.buffer.as_entire_binding(),
                 },
             ],
-            label: Some("ui rsc"),
+            label: Some("ui masks"),
         })
     }
 
     pub fn view_count(&self) -> usize {
         self.textures.view_count()
     }
-}
 
-pub struct UiLimits {
-    max_textures: u32,
-    max_samplers: u32,
-}
+    /// Standalone-image bind groups built since the last call -- see
+    /// `GpuTextures::take_bind_group_creates`. Call once per frame before
+    /// `update()` to measure exactly that frame.
+    pub fn take_image_bind_group_creates(&mut self) -> u64 {
+        self.textures.take_bind_group_creates()
+    }
 
-impl Default for UiLimits {
-    fn default() -> Self {
-        Self {
-            max_textures: 100000,
-            max_samplers: 1000,
-        }
+    /// Atlas-array `grow_array` calls since the last call -- same calling
+    /// convention as `take_image_bind_group_creates` (call once per frame,
+    /// before `update()`, to read exactly the previous frame's tally). Part
+    /// of the Diagnostics page's per-frame report (RUST.md's P0 box, "the
+    /// first input frame" investigation): if a report ever shows a grow
+    /// landing on the same frame the glyphs vanished, that is the
+    /// coincidence to chase first.
+    pub fn take_atlas_pages_grown(&mut self) -> u64 {
+        self.textures.take_pages_grown()
     }
 }
 
-impl UiLimits {
-    pub fn max_binding_array_elements_per_shader_stage(&self) -> u32 {
-        self.max_textures + self.max_samplers
-    }
-    pub fn max_binding_array_sampler_elements_per_shader_stage(&self) -> u32 {
-        self.max_samplers
-    }
+/// What `UiRenderNode::update` changed this frame that a caller building a
+/// per-frame diagnostic report cares about -- see `take_image_bind_group_creates`/
+/// `take_atlas_pages_grown` for the two counters this doesn't carry (they
+/// use the existing "call before update()" convention instead, so as not
+/// to disturb `bench_images`' documented counts).
+#[derive(Clone, Copy, Debug, Default)]
+pub struct FrameUpdateStats {
+    pub masks_resized: bool,
+    pub moves_resized: bool,
 }

core/src/render/primitive.rs

@@ -4,35 +4,27 @@ use crate::{
     Color, UiRegion, WidgetId,
     render::{
         ArrBuf,
-        data::{MaskIdx, PrimitiveInstance},
+        data::{MaskIdx, MoveIdx, PrimitiveInstance},
     },
+    util::HashSet,
 };
 use bytemuck::Pod;
 use wgpu::*;
 
-pub struct Primitives {
-    instances: Vec<PrimitiveInstance>,
-    assoc: Vec<WidgetId>,
-    data: PrimitiveData,
-    free: Vec<usize>,
-    pub updated: bool,
-}
-
-impl Default for Primitives {
-    fn default() -> Self {
-        Self {
-            instances: Default::default(),
-            assoc: Default::default(),
-            data: Default::default(),
-            free: Vec::new(),
-            updated: true,
-        }
-    }
-}
+/// The `binding` tag `Painter` writes on an image instance. Distinct from any
+/// `Primitive::BINDING` because images have no `PrimitiveData` entry to key
+/// one from -- a bind group already selects the texture -- so this only ever
+/// has to match the shader's `TEXTURE` constant and flag "this instance is
+/// drawn with its own bind group" to the code below.
+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 {
@@ -54,6 +46,14 @@ macro_rules! primitives {
 
         impl PrimitiveBuffers {
             pub const LEN: usize = primitives!(@count $($name)*);
+            /// The group-1 binding number each primitive's storage buffer
+            /// sits at, in declaration order. Not `0..LEN`: a primitive's
+            /// `BINDING` also tags its instances for the shader's dispatch
+            /// switch, and a removed primitive (as `TEXTURE` was, once
+            /// images stopped needing a per-instance storage entry) can
+            /// leave a gap, so the pipeline layout has to ask for these
+            /// exact numbers rather than assuming they are contiguous.
+            pub const BINDINGS: [u32; Self::LEN] = [$(<$ty>::BINDING,)*];
             pub fn buffers(&self) -> [(u32, &Buffer); Self::LEN] {
                 [
                     $((<$ty>::BINDING, &self.$name.buffer),)*
@@ -90,73 +90,236 @@ macro_rules! primitives {
                 fn vec(data: &mut PrimitiveData) -> &mut PrimitiveVec<Self> {
                     &mut data.$name
                 }
+                fn vec_ref(data: &PrimitiveData) -> &PrimitiveVec<Self> {
+                    &data.$name
+                }
             }
         )*
     };
-    (@count $t1:tt $($t:tt)+) => { 1 + primitives!(@count $($t),+) };
+    // The recursion has to hand back the same shape it matches -- space
+    // separated, not comma separated. Written with `$($t),+` it re-entered
+    // with a comma as the first token and never terminated, which happened to
+    // work only because there were exactly two primitives: the first step left
+    // a single token, and a single token matches the base case whichever
+    // separator it was written with.
+    (@count $t1:tt $($t:tt)+) => { 1 + primitives!(@count $($t)+) };
     (@count $t:tt) => { 1 };
 }
 
-pub struct PrimitiveInst<P> {
-    pub id: WidgetId,
-    pub primitive: P,
-    pub region: UiRegion,
-    pub mask_idx: MaskIdx,
+/// Every primitive instance in the tree, in one arena that all layers
+/// share, plus the per-primitive data (`rects`, `glyphs`) they index.
+///
+/// **Why one arena rather than one per layer**, which is what this was:
+/// the fragment stage evaluates a *mask's* primitive at the masked pixel
+/// (LAYOUT.md's "Masks with a shape"), and the widget that owns a mask is
+/// routinely in a different layer from the content it clips -- a rounded
+/// container in one layer, a `Stack`'s child content in the layer below.
+/// A per-layer buffer cannot answer that lookup at all: only one layer's
+/// group is bound at a time, so the mask would silently read another
+/// layer's rect. Both buffers are therefore global and bound once per
+/// frame, and a layer keeps only its draw *order* ([`LayerOrder`]).
+///
+/// Slots are stable for a primitive's whole life: nothing here is
+/// compacted, so a `Mask` can hold a slot across frames.
+pub struct Primitives {
+    instances: Vec<PrimitiveInstance>,
+    assoc: Vec<WidgetId>,
+    /// Where each slot's [`PrimitiveHandle`] sits in its owner's
+    /// `ActiveData::primitives` -- the index that makes
+    /// `UiRenderState::apply_free` O(1) per renumbered primitive instead
+    /// of a scan of everything the owner drew. Written by
+    /// [`Self::set_handle_index`] from the one place a handle is taken
+    /// into that vec (`Painter::own`), and dead alongside its `assoc`
+    /// entry, which is what keeps the two in step.
+    ///
+    /// Without it a text widget that is freed and redrawn in one frame
+    /// costs O(glyphs^2): every one of its glyphs is renumbered, and each
+    /// renumbering scanned all of them. Measured 2026-09-08 at 1.37s for a
+    /// 51,200-glyph block on this machine, against 20ms for the shaping
+    /// and rasterising of the same text.
+    handle_idx: Vec<u32>,
+    /// Slots freed since the last [`Self::apply_free`]. Deliberately not
+    /// reusable yet: the layer that drew one still names it in its draw
+    /// order until that call compacts the order, so handing it out again
+    /// first would draw the new primitive twice -- once through the stale
+    /// order entry and once through the new one.
+    freed: Vec<usize>,
+    /// Slots [`Self::apply_free`] released, which is what [`Self::alloc`]
+    /// hands out.
+    reusable: Vec<usize>,
+    data: PrimitiveData,
+    /// Whether the instance arena or the per-primitive data changed since
+    /// the last upload -- one flag for both, since they are uploaded
+    /// together.
+    pub updated: bool,
+}
+
+impl Default for Primitives {
+    fn default() -> Self {
+        Self {
+            instances: Default::default(),
+            assoc: Default::default(),
+            handle_idx: Default::default(),
+            freed: Vec::new(),
+            reusable: Vec::new(),
+            data: Default::default(),
+            updated: true,
+        }
+    }
 }
 
 impl Primitives {
-    pub fn write<P: Primitive>(
+    /// A slot whose handle has not been recorded yet -- see
+    /// [`Self::handle_idx`]. No owner draws four billion primitives, so
+    /// the sentinel cannot collide with a real index.
+    const NO_HANDLE: u32 = u32::MAX;
+
+    /// Writes a primitive into the arena and hands back its slot and its
+    /// entry in the per-primitive data. The caller (`UiRenderState`) puts
+    /// the slot into a layer's draw order -- an instance that no layer
+    /// names is never rasterized, which is what a mask shape drawn only to
+    /// be *referenced* uses.
+    pub fn alloc<P: Primitive>(
         &mut self,
-        layer: usize,
         PrimitiveInst {
             id,
             primitive,
             region,
             mask_idx,
+            move_idx,
         }: PrimitiveInst<P>,
-    ) -> PrimitiveHandle {
+    ) -> (u32, usize) {
+        let data_idx = P::vec(&mut self.data).add(primitive);
+        let slot = self.push(
+            PrimitiveInstance {
+                region,
+                idx: data_idx as u32,
+                mask_idx,
+                move_idx,
+                binding: P::BINDING,
+            },
+            id,
+        );
+        (slot, data_idx)
+    }
+
+    /// A standalone image, which has no `PrimitiveData` entry to allocate
+    /// -- its bind group already picks the texture, so `texture_idx` rides
+    /// in the otherwise-unused `idx` field and names the bind group the
+    /// draw call selects.
+    pub fn alloc_image(
+        &mut self,
+        id: WidgetId,
+        texture_idx: u32,
+        region: UiRegion,
+        mask_idx: MaskIdx,
+        move_idx: MoveIdx,
+    ) -> u32 {
+        self.push(
+            PrimitiveInstance {
+                region,
+                idx: texture_idx,
+                mask_idx,
+                move_idx,
+                binding: IMAGE_BINDING,
+            },
+            id,
+        )
+    }
+
+    fn push(&mut self, inst: PrimitiveInstance, id: WidgetId) -> u32 {
         self.updated = true;
-        let vec = P::vec(&mut self.data);
-        let i = vec.add(primitive);
-        let inst = PrimitiveInstance {
-            region,
-            idx: i as u32,
-            mask_idx,
-            binding: P::BINDING,
-        };
-        let inst_i = if let Some(i) = self.free.pop() {
+        let slot = if let Some(i) = self.reusable.pop() {
             self.instances[i] = inst;
             self.assoc[i] = id;
+            self.handle_idx[i] = Self::NO_HANDLE;
             i
         } else {
-            let i = self.instances.len();
             self.instances.push(inst);
             self.assoc.push(id);
-            i
+            self.handle_idx.push(Self::NO_HANDLE);
+            self.instances.len() - 1
         };
-        PrimitiveHandle::new::<P>(layer, inst_i, i)
-    }
-
-    /// returns (old index, new index)
-    pub fn apply_free(&mut self) -> impl Iterator<Item = PrimitiveChange> {
-        self.free.sort_by(|a, b| b.cmp(a));
-        self.free.drain(..).filter_map(|i| {
-            self.instances.swap_remove(i);
-            self.assoc.swap_remove(i);
-            if i == self.instances.len() {
-                return None;
-            }
-            let id = self.assoc[i];
-            let old = self.instances.len();
-            Some(PrimitiveChange { id, old, new: i })
-        })
+        slot as u32
     }
 
+    /// Retires a slot, answering the mask it was drawn under so the caller
+    /// can drop that mask's ref. The slot itself only becomes reusable at
+    /// the next [`Self::apply_free`] -- see `freed`.
     pub fn free(&mut self, h: &PrimitiveHandle) -> MaskIdx {
         self.updated = true;
-        self.data.free(h.binding, h.data_idx);
-        self.free.push(h.inst_idx);
-        self.instances[h.inst_idx].mask_idx
+        let slot = h.slot as usize;
+        if h.binding != IMAGE_BINDING {
+            self.data.free(h.binding, h.data_idx);
+        }
+        self.freed.push(slot);
+        self.instances[slot].mask_idx
+    }
+
+    /// Hands this frame's freed slots back for reuse. Called once per
+    /// frame from `UiRenderState::update`, **after** every layer has
+    /// compacted its draw order, since that order is the only thing still
+    /// naming them.
+    pub fn release_freed(&mut self) {
+        self.reusable.append(&mut self.freed);
+    }
+
+    /// Which widget drew the primitive in `slot` -- how a draw-order
+    /// change finds the handle it has to renumber.
+    pub fn owner(&self, slot: u32) -> WidgetId {
+        self.assoc[slot as usize]
+    }
+
+    /// Records that `slot`'s handle is `idx` entries into its owner's
+    /// `ActiveData::primitives`. Called once per primitive, by the one
+    /// place that puts a handle into that vec.
+    pub fn set_handle_index(&mut self, slot: u32, idx: u32) {
+        self.handle_idx[slot as usize] = idx;
+    }
+
+    /// Where `slot`'s handle sits in its owner's `ActiveData::primitives`
+    /// -- see [`Self::handle_idx`]. `None` only for a slot whose owner
+    /// never took the handle, which nothing in this crate does.
+    pub fn handle_index(&self, slot: u32) -> Option<usize> {
+        match self.handle_idx[slot as usize] {
+            Self::NO_HANDLE => None,
+            idx => Some(idx as usize),
+        }
+    }
+
+    pub fn clear(&mut self) {
+        self.updated = true;
+        self.instances.clear();
+        self.assoc.clear();
+        self.handle_idx.clear();
+        self.freed.clear();
+        self.reusable.clear();
+        self.data.clear();
+    }
+
+    /// How many instances are still live -- the O(1) half of the orphan
+    /// check, so the O(primitives) walk below only runs on a frame that
+    /// already looks wrong. See
+    /// [`crate::UiRenderState::orphaned_primitives`].
+    pub fn live_count(&self) -> usize {
+        self.instances.len() - self.freed.len() - self.reusable.len()
+    }
+
+    /// Every live instance as `(slot, owner, is_image)` -- everything
+    /// except the freed and the reusable. Only
+    /// [`crate::UiRenderState::orphaned_primitives`] uses this, to check
+    /// that every live primitive still belongs to a live widget.
+    pub fn live_instances(&self) -> impl Iterator<Item = (u32, WidgetId, bool)> + '_ {
+        let dead: HashSet<usize> = self.freed.iter().chain(&self.reusable).copied().collect();
+        (0..self.instances.len())
+            .filter(move |i| !dead.contains(i))
+            .map(|i| {
+                (
+                    i as u32,
+                    self.assoc[i],
+                    self.instances[i].binding == IMAGE_BINDING,
+                )
+            })
     }
 
     pub fn data(&self) -> &PrimitiveData {
@@ -167,40 +330,166 @@ impl Primitives {
         &self.instances
     }
 
+    pub fn instance(&self, slot: u32) -> &PrimitiveInstance {
+        &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.inst_idx].region
+        &mut self.instances[h.slot as usize].region
     }
 }
 
-pub struct PrimitiveChange {
-    pub id: WidgetId,
-    pub old: usize,
-    pub new: usize,
+/// One layer's draw order: the slots of the global arena it draws, in the
+/// order they were written. The vertex buffer of a layer is exactly this.
+///
+/// Both lists free with `swap_remove`, so a layer's draw order was already
+/// undefined before this split: nothing here may assume one primitive
+/// stays adjacent to another once anything in the layer has been freed.
+#[derive(Default)]
+pub struct LayerOrder {
+    order: Vec<u32>,
+    /// Standalone images, kept apart because each draws with its own bind
+    /// group rather than sharing the layer's one instanced draw -- see
+    /// `UiRenderNode::draw`.
+    images: Vec<u32>,
+    free: Vec<usize>,
+    image_free: Vec<usize>,
+    pub updated: bool,
 }
 
+impl LayerOrder {
+    pub fn push(&mut self, slot: u32, is_image: bool) -> usize {
+        self.updated = true;
+        let list = if is_image {
+            &mut self.images
+        } else {
+            &mut self.order
+        };
+        list.push(slot);
+        list.len() - 1
+    }
+
+    /// Marks a position for removal. Deferred to [`Self::apply_free`] like
+    /// the arena's own, so that a position is only renumbered once per
+    /// frame however many were dropped.
+    pub fn free(&mut self, pos: usize, is_image: bool) {
+        self.updated = true;
+        if is_image {
+            self.image_free.push(pos);
+        } else {
+            self.free.push(pos);
+        }
+    }
+
+    /// Compacts both lists, answering every primitive whose position
+    /// moved so its handle can be corrected.
+    pub fn apply_free(&mut self) -> Vec<OrderChange> {
+        let mut changes = Self::apply_free_list(&mut self.free, &mut self.order, false);
+        changes.extend(Self::apply_free_list(
+            &mut self.image_free,
+            &mut self.images,
+            true,
+        ));
+        changes
+    }
+
+    fn apply_free_list(
+        free: &mut Vec<usize>,
+        list: &mut Vec<u32>,
+        is_image: bool,
+    ) -> Vec<OrderChange> {
+        // Descending, so removing a contiguous tail costs no renumbering
+        // at all -- which is what freeing one widget's primitives is.
+        free.sort_by(|a, b| b.cmp(a));
+        free.drain(..)
+            .filter_map(|pos| {
+                list.swap_remove(pos);
+                if pos == list.len() {
+                    return None;
+                }
+                Some(OrderChange {
+                    slot: list[pos],
+                    is_image,
+                    pos,
+                })
+            })
+            .collect()
+    }
+
+    pub fn order(&self) -> &Vec<u32> {
+        &self.order
+    }
+
+    pub fn images(&self) -> &Vec<u32> {
+        &self.images
+    }
+}
+
+/// A primitive whose position in a layer's draw order moved when
+/// something before it was freed -- `slot` names which primitive, so its
+/// owner's handle can be found and pointed at `pos`.
+pub struct OrderChange {
+    pub slot: u32,
+    /// Which of the layer's two lists moved: their positions are
+    /// independent index spaces, so a handle matching on position alone
+    /// could take an image's renumbering for a rect's.
+    pub is_image: bool,
+    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 ([`NOT_DRAWN`] if it is
+/// only referenced).
 #[derive(Debug)]
 pub struct PrimitiveHandle {
     pub layer: usize,
-    pub inst_idx: usize,
+    pub pos: usize,
+    pub slot: u32,
     pub data_idx: usize,
     pub binding: u32,
 }
 
 impl PrimitiveHandle {
-    fn new<P: Primitive>(layer: usize, inst_idx: usize, data_idx: usize) -> Self {
-        Self {
-            layer,
-            inst_idx,
-            data_idx,
-            binding: P::BINDING,
-        }
+    pub fn is_image(&self) -> bool {
+        self.binding == IMAGE_BINDING
     }
 }
 
+pub struct PrimitiveInst<P> {
+    pub id: WidgetId,
+    pub primitive: P,
+    pub region: UiRegion,
+    pub mask_idx: MaskIdx,
+    pub move_idx: MoveIdx,
+}
+
 primitives!(
     rects: RectPrimitive => 0,
-    textures: TexturePrimitive => 1,
+    glyphs: GlyphPrimitive => 2,
 );
 
 #[repr(C)]
@@ -223,11 +512,48 @@ impl RectPrimitive {
     }
 }
 
+/// One glyph, drawn as a sub-rectangle of the glyph atlas array.
+///
+/// `color` is the text colour and is multiplied by the atlas's alpha for an
+/// ordinary mask glyph; a colour glyph (emoji) carries its own colour and
+/// takes the atlas texel unchanged, which is what `IS_COLOR` selects.
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
-pub struct TexturePrimitive {
-    pub view_idx: u32,
-    pub sampler_idx: u32,
+pub struct GlyphPrimitive {
+    pub uv_min: [f32; 2],
+    pub uv_max: [f32; 2],
+    /// Layer of the shared atlas array texture this glyph's page occupies --
+    /// not a bind-group or view index, since a page never gets one of its
+    /// own. See TEXTURES.md's "Recommended shape".
+    pub layer: u32,
+    pub color: Color<u8>,
+    pub flags: u32,
+    /// Pads this struct's Rust size to match WGSL's storage-buffer layout for
+    /// `GlyphInfo`: two `vec2<f32>` members give the struct an 8-byte
+    /// alignment, which rounds the WGSL size up to 32 bytes even though the
+    /// fields above only total 28. `bytemuck` does not check this for us.
+    _pad: u32,
+}
+
+impl GlyphPrimitive {
+    pub const IS_COLOR: u32 = 1;
+
+    pub fn new(
+        uv_min: [f32; 2],
+        uv_max: [f32; 2],
+        layer: u32,
+        color: Color<u8>,
+        flags: u32,
+    ) -> Self {
+        Self {
+            uv_min,
+            uv_max,
+            layer,
+            color,
+            flags,
+            _pad: 0,
+        }
+    }
 }
 
 pub struct PrimitiveVec<T> {

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)
+}

core/src/render/shader.wgsl

@@ -1,12 +1,16 @@
 const RECT: u32 = 0u;
+// TEXTURE has no entry in group 1: a standalone image draws with its own
+// bind group (see UiRenderNode::draw), so there is nothing per-instance left
+// to look up here -- the bind group already picked the texture.
 const TEXTURE: u32 = 1u;
+const GLYPH: u32 = 2u;
 
 @group(0) @binding(0)
 var<uniform> window: WindowUniform;
 @group(1) @binding(RECT)
 var<storage> rects: array<Rect>;
-@group(1) @binding(TEXTURE)
-var<storage> textures: array<TextureInfo>;
+@group(1) @binding(GLYPH)
+var<storage> glyphs: array<GlyphInfo>;
 
 struct Rect {
     color: u32,
@@ -15,14 +19,30 @@ struct Rect {
     inner_radius: f32,
 }
 
-struct TextureInfo {
-    view_idx: u32,
-    sampler_idx: u32,
+struct GlyphInfo {
+    uv_min: vec2<f32>,
+    uv_max: vec2<f32>,
+    // Layer of the shared atlas array texture, not a view or bind-group
+    // index -- a page never gets its own bind group. See TEXTURES.md's
+    // "Recommended shape".
+    layer: u32,
+    color: u32,
+    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,
+    primitive: u32,
+    parent: u32,
+}
+
+/// One widget's cumulative on-screen translation and the slot of the
+/// ancestor to add on top of it. Mirrors `MoveOffset` in data.rs.
+struct MoveOffset {
+    delta: vec2<f32>,
+    parent: u32,
 }
 
 struct UiSpan {
@@ -35,39 +55,93 @@ 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>>
+// this replaced, which needed VK_EXT_descriptor_indexing and does not survive
+// a real share of Android GPUs (see TEXTURES.md).
 @group(2) @binding(0)
-var views: binding_array<texture_2d<f32>>;
+var atlas: texture_2d_array<f32>;
+// One standalone image's texture. The main draw (rects and glyphs) binds a
+// 1x1 null texture here, since neither samples it; each image draw call
+// binds its own -- see UiRenderNode::draw.
 @group(2) @binding(1)
-var samplers: binding_array<sampler>;
+var image_texture: texture_2d<f32>;
 @group(2) @binding(2)
+var samp: sampler;
+// Their own group, bound once per frame rather than folded into group 2: see
+// UiRenderNode::masks_layout for why an image's own bind group must not name
+// either buffer.
+@group(3) @binding(0)
 var<storage> masks: array<Mask>;
+@group(3) @binding(1)
+var<storage> move_offsets: array<MoveOffset>;
+// Every primitive's placement, in one arena all layers share. The vertex
+// stage reads the primitive it is drawing (its slot arrives as the only
+// vertex attribute); the fragment stage reads a *mask's* primitive, which
+// is generally a different one in a different layer. See LAYOUT.md's
+// "Masks with a shape" and `Primitives` in primitive.rs.
+@group(3) @binding(2)
+var<storage> instances: array<PrimitiveInstance>;
+
+// The bound on the parent walk, kept in step with `PARENT_CHAIN_LIMIT` in
+// render_state.rs, which walks the identical chain on the CPU side for
+// hit-testing. Bounded so a malformed chain (a cyclic `parent`) cannot
+// hang the GPU -- not a claim about how deep a real tree gets. It was 16
+// and that was too small: the transcript screen's composer field sits 17
+// slots below the root, measured 2026-09-07 on this checkout's emulator
+// by tapping it (the CPU walk's own debug assert names the chain now).
+// Past the bound both walks simply stop summing, so the widget draws and
+// hit-tests short by whatever the outer slots held, with nothing on
+// screen to say so.
+const PARENT_CHAIN_LIMIT: u32 = 64u;
+
+/// Sums the pixel delta along the parent chain starting at `idx`, shared by
+/// the vertex stage (a primitive's own corners) and the fragment stage (its
+/// mask's corners) so the walk is written once. See LAYOUT.md section 2b.
+fn resolve_move(idx: u32) -> vec2<f32> {
+    var total = vec2<f32>(0.0, 0.0);
+    var i = idx;
+    for (var step = 0u; step < PARENT_CHAIN_LIMIT; step++) {
+        let entry = move_offsets[i];
+        total += entry.delta;
+        if entry.parent == 4294967295u {
+            break;
+        }
+        i = entry.parent;
+    }
+    return total;
+}
 
 struct WindowUniform {
     dim: vec2<f32>,
 };
 
+/// Mirrors `PrimitiveInstance` in data.rs -- the placement and what to
+/// draw there. `x`/`y` are the `UiRegion`'s two spans.
+struct PrimitiveInstance {
+    x: UiSpan,
+    y: UiSpan,
+    binding: u32,
+    idx: u32,
+    mask_idx: u32,
+    move_idx: u32,
+}
+
+/// A layer's draw order: one slot into `instances` per instance drawn.
 struct InstanceInput {
-    @location(0) x_start: vec2<f32>,
-    @location(1) x_end: vec2<f32>,
-    @location(2) y_start: vec2<f32>,
-    @location(3) y_end: vec2<f32>,
-    @location(4) binding: u32,
-    @location(5) idx: u32,
-    @location(6) mask_idx: u32,
+    @location(0) slot: u32,
 }
 
 struct VertexOutput {
     @location(0) top_left: vec2<f32>,
     @location(1) bot_right: vec2<f32>,
     @location(2) uv: vec2<f32>,
-    @location(3) binding: u32,
-    @location(4) idx: u32,
-    @location(5) mask_idx: u32,
+    // `flat` is the only interpolation an integer can have, and naga
+    // (wgpu 30) now requires saying so rather than inferring it.
+    @location(3) @interpolate(flat) binding: u32,
+    @location(4) @interpolate(flat) idx: u32,
+    @location(5) @interpolate(flat) mask_idx: u32,
     @builtin(position) clip_position: vec4<f32>,
 };
 
@@ -78,20 +152,38 @@ 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,
     in: InstanceInput,
 ) -> VertexOutput {
     var out: VertexOutput;
+    let inst = instances[in.slot];
 
-    let top_left_rel = vec2(in.x_start.x, in.y_start.x);
-    let top_left_abs = vec2(in.x_start.y, in.y_start.y);
-    let bot_right_rel = vec2(in.x_end.x, in.y_end.x);
-    let bot_right_abs = vec2(in.x_end.y, in.y_end.y);
-
-    let top_left = floor(top_left_rel * window.dim) + floor(top_left_abs);
-    let bot_right = floor(bot_right_rel * window.dim) + floor(bot_right_abs);
+    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<f32>(
@@ -101,11 +193,11 @@ fn vs_main(
     let pos = (top_left + uv * size) / window.dim * 2.0 - 1.0;
     out.clip_position = vec4<f32>(pos.x, -pos.y, 0.0, 1.0);
     out.uv = uv;
-    out.binding = in.binding;
-    out.idx = in.idx;
+    out.binding = inst.binding;
+    out.idx = inst.idx;
     out.top_left = top_left;
     out.bot_right = bot_right;
-    out.mask_idx = in.mask_idx;
+    out.mask_idx = inst.mask_idx;
 
     return out;
 }
@@ -123,29 +215,79 @@ fn fs_main(
             color = draw_rounded_rect(region, rects[i]);
         }
         case TEXTURE: {
-            color = draw_texture(region, textures[i]);
+            color = draw_texture(region);
+        }
+        case GLYPH: {
+            color = draw_glyph(region, glyphs[i]);
         }
         default: {
             color = vec4(1.0, 0.0, 1.0, 1.0);
         }
     }
-    if in.mask_idx != 4294967295u {
-        let mask = masks[in.mask_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);
-        let bot_right = floor(br.rel * window.dim) + floor(br.abs);
-        if pos.x < top_left.x || pos.x > bot_right.x || pos.y < top_left.y || pos.y > bot_right.y {
-            color *= 0.0;
+    // Every mask on the chain, not just the innermost: a widget that set
+    // 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];
+        color.a *= mask_coverage(pos, mask);
+        mask_idx = mask.parent;
     }
     return color;
 }
 
-// TODO: this seems really inefficient (per frag indexing)?
-fn draw_texture(region: Region, info: TextureInfo) -> vec4<f32> {
-    return textureSample(views[info.view_idx], samplers[info.sampler_idx], region.uv);
+/// 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);
+}
+
+fn draw_glyph(region: Region, g: GlyphInfo) -> vec4<f32> {
+    let uv = mix(g.uv_min, g.uv_max, region.uv);
+    let texel = textureSample(atlas, samp, uv, i32(g.layer));
+    if (g.flags & 1u) != 0u {
+        return texel;
+    }
+    var color = unpack4x8unorm(g.color);
+    color.a *= texel.a;
+    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> {
@@ -153,14 +295,12 @@ fn draw_rounded_rect(region: Region, rect: Rect) -> vec4<f32> {
 
     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);
     }

core/src/render/texture.rs

@@ -1,59 +1,306 @@
-use image::{DynamicImage, EncodableLayout};
+use image::{DynamicImage, EncodableLayout, GenericImageView};
 use wgpu::{util::DeviceExt, *};
 
-use crate::{TextureUpdate, Textures};
+use crate::{PatchRect, TextureKind, TextureUpdate, Textures};
 
+use super::atlas::PAGE;
+
+/// The fewest layers the glyph atlas array is ever created with. Two, not
+/// one, for the GLES reason written on `create_array_texture`.
+const MIN_ARRAY_LAYERS: u32 = 2;
+
+/// What one texture slot is, GPU-side. Parallel to `Textures`' own slot
+/// numbering (`TextureKind`'s `Image`/`Page`), so a slot's index means the
+/// same thing on both sides without a second map to keep in sync.
+enum Slot {
+    /// A slot that was freed, or pushed and freed within the same batch
+    /// before ever reaching here.
+    Empty,
+    Image(ImageGpu),
+    /// The array layer a page occupies. Pages are never freed (see
+    /// `Textures::free`), so this is the only variant that outlives a `Free`.
+    Page(u32),
+}
+
+struct ImageGpu {
+    /// Kept alive alongside `view`/`bind_group`, which borrow from it only in
+    /// the sense that dropping this drops the GPU resource they point to.
+    #[allow(dead_code)]
+    texture: Texture,
+    view: TextureView,
+    bind_group: BindGroup,
+}
+
+/// Owns the two kinds of texture iris draws:
+///
+/// - **The glyph atlas**, one `texture_2d_array` whose layers are pages
+///   (`Slot::Page`), grown by recreating the array with headroom and
+///   `copy_texture_to_texture`-ing the old layers across. No feature beyond
+///   Vulkan 1.0/GLES sampling is needed for this -- a layer index is an
+///   ordinary sampling operand.
+/// - **Standalone images** (`Slot::Image`), each its own `Texture` and
+///   `BindGroup`, drawn one `draw()` call at a time with that bind group
+///   bound -- see `UiRenderNode::draw`.
+///
+/// See TEXTURES.md's "Recommended shape" for why, and RUST.md's
+/// "iris's binding array does not survive real Android hardware" for what
+/// this replaced (one giant `binding_array<texture_2d<f32>>` needing
+/// `VK_EXT_descriptor_indexing`, which a real share of Android GPUs lack).
 pub struct GpuTextures {
     device: Device,
     queue: Queue,
-    views: Vec<TextureView>,
-    view_count: usize,
-    samplers: Vec<Sampler>,
+
+    slots: Vec<Slot>,
+
+    array_texture: Texture,
+    array_view: TextureView,
+    array_capacity: u32,
+    /// Layers actually written. Only grows -- see `Slot::Page`.
+    page_count: u32,
+
+    sampler: Sampler,
+    /// Bound in the image slot of the main draw's bind group, which has
+    /// nothing of its own to put there: rects and glyphs never sample it,
+    /// but the layout requires something bound regardless.
     null_view: TextureView,
-    no_views: Vec<TextureView>,
+
+    /// Standalone-image bind groups actually built (`create_image`'s own
+    /// build, or one per slot touched by `rebuild_image_bind_groups`) since
+    /// the last `take_bind_group_creates`. IRIS_TODO.md's "many images"
+    /// benchmark reads this to prove the steady-state cost of an
+    /// unchanging image list is zero, the same way `UiRenderState`'s
+    /// `draw_count`/`region_mut_count` prove the layout side.
+    bind_group_creates: u64,
+    /// `grow_array` calls since the last `take_pages_grown` -- the
+    /// Diagnostics page's per-frame report (RUST.md's P0 box, "the first
+    /// input frame" investigation) reads this alongside `bind_group_creates`
+    /// to say whether *this* frame's glyph disappearance, if any, coincided
+    /// with the atlas array being recreated.
+    pages_grown: u64,
 }
 
 impl GpuTextures {
-    pub fn update(&mut self, textures: &mut Textures) -> bool {
-        let mut changed = false;
+    /// Applies queued `Textures` updates, then reports whether the *main*
+    /// bind group (the one rects and glyphs draw with) needs rebuilding --
+    /// true exactly when the atlas array was recreated (its view identity
+    /// changed). Pushing or freeing a standalone image never touches that
+    /// group: it built or drops its own. Masks/move_offsets resizing is
+    /// `UiRenderNode`'s own concern now (its `masks_group`, group 3) --
+    /// see that struct's field comment for why standalone images no longer
+    /// hear about either buffer at all.
+    pub fn update(&mut self, textures: &mut Textures, rsc_layout: &BindGroupLayout) -> bool {
+        let mut rebuild_main = false;
         for update in textures.updates() {
-            changed = true;
             match update {
-                TextureUpdate::Push(image) => self.push(image),
-                TextureUpdate::Set(i, image) => self.set(i, image),
-                TextureUpdate::SetFree => self.view_count += 1,
+                TextureUpdate::Push(kind, image) => {
+                    rebuild_main |= self.push(kind, image, rsc_layout);
+                }
+                TextureUpdate::Set(kind, i, image) => {
+                    rebuild_main |= self.set(kind, i, image, rsc_layout);
+                }
+                // A patch changes texture contents, not which layer or bind
+                // group exists, so it never asks for a rebuild -- rebuilding
+                // per glyph is exactly the cost this exists to avoid.
+                TextureUpdate::Patch(i, rect, image) => self.patch(i, rect, image),
+                TextureUpdate::SetFree => {}
                 TextureUpdate::Free(i) => self.free(i),
-                TextureUpdate::PushFree => self.push_free(),
+                TextureUpdate::PushFree(_kind) => self.slots.push(Slot::Empty),
             }
         }
-        changed
-    }
-    fn set(&mut self, i: u32, image: &DynamicImage) {
-        self.view_count += 1;
-        let view = self.create_view(image);
-        self.views[i as usize] = view;
-    }
-    fn free(&mut self, i: u32) {
-        self.view_count -= 1;
-        self.views[i as usize] = self.null_view.clone();
-    }
-    fn push(&mut self, image: &DynamicImage) {
-        self.view_count += 1;
-        let view = self.create_view(image);
-        self.views.push(view);
-    }
-    fn push_free(&mut self) {
-        self.view_count += 1;
-        self.views.push(self.null_view.clone());
+        rebuild_main
     }
 
-    fn create_view(&self, image: &DynamicImage) -> TextureView {
-        let image = image.to_rgba8();
-        let (width, height) = image.dimensions();
+    fn push(
+        &mut self,
+        kind: TextureKind,
+        image: &DynamicImage,
+        rsc_layout: &BindGroupLayout,
+    ) -> bool {
+        let (slot, rebuilt) = self.make_slot(kind, image, rsc_layout);
+        self.slots.push(slot);
+        rebuilt
+    }
+
+    fn set(
+        &mut self,
+        kind: TextureKind,
+        i: u32,
+        image: &DynamicImage,
+        rsc_layout: &BindGroupLayout,
+    ) -> bool {
+        let (slot, rebuilt) = self.make_slot(kind, image, rsc_layout);
+        self.slots[i as usize] = slot;
+        rebuilt
+    }
+
+    fn make_slot(
+        &mut self,
+        kind: TextureKind,
+        image: &DynamicImage,
+        rsc_layout: &BindGroupLayout,
+    ) -> (Slot, bool) {
+        match kind {
+            TextureKind::Image => {
+                let gpu = self.create_image(image, rsc_layout);
+                (Slot::Image(gpu), false)
+            }
+            TextureKind::Page { layer } => {
+                let mut rebuilt = false;
+                if layer >= self.array_capacity {
+                    self.grow_array(rsc_layout);
+                    rebuilt = true;
+                }
+                self.write_full_layer(layer, image);
+                self.page_count = self.page_count.max(layer + 1);
+                (Slot::Page(layer), rebuilt)
+            }
+        }
+    }
+
+    fn free(&mut self, i: u32) {
+        if let Some(slot) = self.slots.get_mut(i as usize) {
+            *slot = Slot::Empty;
+        }
+        // A page's layer is not reclaimed here either -- see `Slot::Page`.
+    }
+
+    fn patch(&mut self, i: u32, rect: PatchRect, image: &DynamicImage) {
+        let Some(&Slot::Page(layer)) = self.slots.get(i as usize) else {
+            return;
+        };
+        if rect.width == 0 || rect.height == 0 {
+            return;
+        }
+        // Cropped rather than written straight from the atlas, because
+        // write_texture wants tightly packed rows and the atlas rows are as
+        // wide as the atlas. A glyph is small, so the copy is too.
+        let sub = image
+            .view(rect.x, rect.y, rect.width, rect.height)
+            .to_image();
+        self.queue.write_texture(
+            TexelCopyTextureInfo {
+                texture: &self.array_texture,
+                mip_level: 0,
+                origin: Origin3d {
+                    x: rect.x,
+                    y: rect.y,
+                    z: layer,
+                },
+                aspect: TextureAspect::All,
+            },
+            sub.as_bytes(),
+            TexelCopyBufferLayout {
+                offset: 0,
+                bytes_per_row: Some(rect.width * 4),
+                rows_per_image: Some(rect.height),
+            },
+            Extent3d {
+                width: rect.width,
+                height: rect.height,
+                depth_or_array_layers: 1,
+            },
+        );
+    }
+
+    fn write_full_layer(&self, layer: u32, image: &DynamicImage) {
+        // Every page is created as exactly PAGE x PAGE (`GlyphAtlas::allocate`),
+        // so this is always a whole-layer write, never a crop.
+        let rgba = image.to_rgba8();
+        self.queue.write_texture(
+            TexelCopyTextureInfo {
+                texture: &self.array_texture,
+                mip_level: 0,
+                origin: Origin3d {
+                    x: 0,
+                    y: 0,
+                    z: layer,
+                },
+                aspect: TextureAspect::All,
+            },
+            rgba.as_bytes(),
+            TexelCopyBufferLayout {
+                offset: 0,
+                bytes_per_row: Some(PAGE * 4),
+                rows_per_image: Some(PAGE),
+            },
+            Extent3d {
+                width: PAGE,
+                height: PAGE,
+                depth_or_array_layers: 1,
+            },
+        );
+    }
+
+    /// Doubles the array's layer capacity (headroom, so this is rare) and
+    /// copies the old layers across GPU-side -- no readback. Recreates the
+    /// array's view, which invalidates every bind group that referenced it,
+    /// so this also rebuilds all of them before returning.
+    fn grow_array(&mut self, rsc_layout: &BindGroupLayout) {
+        self.pages_grown += 1;
+        let new_capacity = self.array_capacity * 2;
+        let new_texture = Self::create_array_texture(&self.device, new_capacity);
+        if self.page_count > 0 {
+            let mut encoder = self
+                .device
+                .create_command_encoder(&CommandEncoderDescriptor {
+                    label: Some("atlas array grow"),
+                });
+            encoder.copy_texture_to_texture(
+                TexelCopyTextureInfo {
+                    texture: &self.array_texture,
+                    mip_level: 0,
+                    origin: Origin3d::ZERO,
+                    aspect: TextureAspect::All,
+                },
+                TexelCopyTextureInfo {
+                    texture: &new_texture,
+                    mip_level: 0,
+                    origin: Origin3d::ZERO,
+                    aspect: TextureAspect::All,
+                },
+                Extent3d {
+                    width: PAGE,
+                    height: PAGE,
+                    depth_or_array_layers: self.page_count,
+                },
+            );
+            self.queue.submit(std::iter::once(encoder.finish()));
+        }
+        self.array_texture = new_texture;
+        self.array_view = self.array_texture.create_view(&TextureViewDescriptor {
+            dimension: Some(TextureViewDimension::D2Array),
+            ..Default::default()
+        });
+        self.array_capacity = new_capacity;
+        self.rebuild_image_bind_groups(rsc_layout);
+    }
+
+    /// Called only from `grow_array`: the atlas array's view identity is the
+    /// one thing an image's bind group (group 2) still names that can
+    /// change out from under it. Masks/move_offsets resizing no longer
+    /// reaches here at all -- see `UiRenderNode::masks_group`.
+    fn rebuild_image_bind_groups(&mut self, rsc_layout: &BindGroupLayout) {
+        for slot in &mut self.slots {
+            if let Slot::Image(gpu) = slot {
+                gpu.bind_group = Self::make_image_bind_group(
+                    &self.device,
+                    rsc_layout,
+                    &self.array_view,
+                    &gpu.view,
+                    &self.sampler,
+                );
+                self.bind_group_creates += 1;
+            }
+        }
+    }
+
+    fn create_image(&mut self, image: &DynamicImage, rsc_layout: &BindGroupLayout) -> ImageGpu {
+        let rgba = image.to_rgba8();
+        let (width, height) = rgba.dimensions();
         let texture = self.device.create_texture_with_data(
             &self.queue,
             &TextureDescriptor {
-                label: None,
+                label: Some("image"),
                 size: Extent3d {
                     width,
                     height,
@@ -63,45 +310,173 @@ impl GpuTextures {
                 sample_count: 1,
                 dimension: TextureDimension::D2,
                 format: TextureFormat::Rgba8Unorm,
-                usage: TextureUsages::TEXTURE_BINDING,
+                usage: TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_DST,
                 view_formats: &[],
             },
             wgt::TextureDataOrder::MipMajor,
-            image.as_bytes(),
+            rgba.as_bytes(),
         );
-        texture.create_view(&TextureViewDescriptor::default())
+        let view = texture.create_view(&TextureViewDescriptor::default());
+        let bind_group = Self::make_image_bind_group(
+            &self.device,
+            rsc_layout,
+            &self.array_view,
+            &view,
+            &self.sampler,
+        );
+        self.bind_group_creates += 1;
+        ImageGpu {
+            texture,
+            view,
+            bind_group,
+        }
+    }
+
+    /// Builds group 2 for one standalone image: the shared atlas array, this
+    /// image's own view and the shared sampler -- the same layout the main
+    /// draw uses with a null view in the image slot. Deliberately does not
+    /// touch masks/move_offsets (group 3, `UiRenderNode::masks_group`): see
+    /// that field's comment for why folding them in here was the bug.
+    fn make_image_bind_group(
+        device: &Device,
+        rsc_layout: &BindGroupLayout,
+        array_view: &TextureView,
+        image_view: &TextureView,
+        sampler: &Sampler,
+    ) -> BindGroup {
+        device.create_bind_group(&BindGroupDescriptor {
+            layout: rsc_layout,
+            entries: &[
+                BindGroupEntry {
+                    binding: 0,
+                    resource: BindingResource::TextureView(array_view),
+                },
+                BindGroupEntry {
+                    binding: 1,
+                    resource: BindingResource::TextureView(image_view),
+                },
+                BindGroupEntry {
+                    binding: 2,
+                    resource: BindingResource::Sampler(sampler),
+                },
+            ],
+            label: Some("ui rsc image"),
+        })
+    }
+
+    /// The atlas is sampled as a `texture_2d_array`, and **a one-layer
+    /// array is not one on the GLES backend**: wgpu-hal picks the GL
+    /// texture target from the descriptor alone
+    /// (`gles::Texture::get_info_from_desc`, `(false, 1) => TEXTURE_2D`),
+    /// so a capacity of 1 creates a `GL_TEXTURE_2D` and binds it to the
+    /// shader's `sampler2DArray`. GL then treats that unit as incomplete
+    /// and every `textureSample` returns (0, 0, 0, 1) -- which, through
+    /// `draw_glyph`'s `color.a *= texel.a`, draws every glyph as a solid
+    /// filled box. That was iris's appearance on the emulator's GLES for
+    /// two days (RUST.md, "the emulator cannot draw iris's glyphs"), and
+    /// it is a real defect on any device whose adapter is GL rather than
+    /// Vulkan, not an emulator artifact. So the array never has fewer than
+    /// `MIN_ARRAY_LAYERS` layers; the second layer costs one page of
+    /// texture memory and is used by the next atlas page anyway.
+    fn create_array_texture(device: &Device, capacity: u32) -> Texture {
+        debug_assert!(
+            capacity >= MIN_ARRAY_LAYERS,
+            "glyph atlas array asked for {capacity} layers; fewer than {MIN_ARRAY_LAYERS} is a \
+             GL_TEXTURE_2D on the GLES backend and draws every glyph as a box"
+        );
+        device.create_texture(&TextureDescriptor {
+            label: Some("glyph atlas array"),
+            size: Extent3d {
+                width: PAGE,
+                height: PAGE,
+                depth_or_array_layers: capacity,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: TextureDimension::D2,
+            format: TextureFormat::Rgba8Unorm,
+            usage: TextureUsages::TEXTURE_BINDING
+                | TextureUsages::COPY_DST
+                | TextureUsages::COPY_SRC,
+            view_formats: &[],
+        })
     }
 
     pub fn new(device: &Device, queue: &Queue) -> Self {
+        let sampler = default_sampler(device);
         let null_view = null_texture_view(device);
+        let array_capacity = MIN_ARRAY_LAYERS;
+        let array_texture = Self::create_array_texture(device, array_capacity);
+        let array_view = array_texture.create_view(&TextureViewDescriptor {
+            dimension: Some(TextureViewDimension::D2Array),
+            ..Default::default()
+        });
         Self {
             device: device.clone(),
             queue: queue.clone(),
-            views: Vec::new(),
-            samplers: vec![default_sampler(device)],
-            no_views: vec![null_view.clone()],
+            slots: Vec::new(),
+            array_texture,
+            array_view,
+            array_capacity,
+            page_count: 0,
+            sampler,
             null_view,
-            view_count: 0,
+            bind_group_creates: 0,
+            pages_grown: 0,
         }
     }
 
-    pub fn views(&self) -> Vec<&TextureView> {
-        if self.views.is_empty() {
-            &self.no_views
-        } else {
-            &self.views
-        }
-        .iter()
-        .by_ref()
-        .collect()
+    /// Reads and zeroes the standalone-image bind-group creation counter --
+    /// call once per frame before `update()`, mirroring
+    /// `UiRenderState::take_counters`.
+    pub fn take_bind_group_creates(&mut self) -> u64 {
+        std::mem::take(&mut self.bind_group_creates)
     }
 
-    pub fn samplers(&self) -> Vec<&Sampler> {
-        self.samplers.iter().by_ref().collect()
+    /// Reads and zeroes the atlas-array-grow counter -- see `pages_grown`'s
+    /// field comment.
+    pub fn take_pages_grown(&mut self) -> u64 {
+        std::mem::take(&mut self.pages_grown)
+    }
+
+    pub fn array_view(&self) -> &TextureView {
+        &self.array_view
+    }
+
+    pub fn null_view(&self) -> &TextureView {
+        &self.null_view
+    }
+
+    pub fn sampler(&self) -> &Sampler {
+        &self.sampler
+    }
+
+    /// The bind group a standalone image draws with. Panics if `idx` names an
+    /// atlas page or a freed slot instead -- either is a caller bug (the
+    /// wrong kind of instance reached this draw path), not a condition to
+    /// recover from.
+    pub fn image_bind_group(&self, idx: u32) -> &BindGroup {
+        match self.slots.get(idx as usize) {
+            Some(Slot::Image(gpu)) => &gpu.bind_group,
+            other => panic!("texture slot {idx} is not a live standalone image: {other:?}"),
+        }
     }
 
     pub fn view_count(&self) -> usize {
-        self.view_count
+        self.slots
+            .iter()
+            .filter(|s| !matches!(s, Slot::Empty))
+            .count()
+    }
+}
+
+impl std::fmt::Debug for Slot {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        match self {
+            Slot::Empty => write!(f, "Empty"),
+            Slot::Image(_) => write!(f, "Image"),
+            Slot::Page(layer) => write!(f, "Page(layer={layer})"),
+        }
     }
 }
 

core/src/render/util/mod.rs

@@ -21,13 +21,18 @@ impl<T: Pod> ArrBuf<T> {
             _pd: PhantomData,
         }
     }
-    pub fn update(&mut self, device: &Device, queue: &Queue, data: &[T]) {
-        if self.len != data.len() {
+    /// Returns whether the underlying `Buffer` was recreated -- a caller that
+    /// cached a `BindGroup` referencing it (as `GpuTextures` does for the
+    /// masks buffer) needs to know to rebuild that too.
+    pub fn update(&mut self, device: &Device, queue: &Queue, data: &[T]) -> bool {
+        let resized = self.len != data.len();
+        if resized {
             self.len = data.len();
             self.buffer =
                 Self::init_buf(device, std::mem::size_of_val(data), self.usage, self.label);
         }
         queue.write_buffer(&self.buffer, 0, bytemuck::cast_slice(data));
+        resized
     }
     fn init_buf(device: &Device, size: usize, usage: BufferUsages, label: &'static str) -> Buffer {
         let mut size = size as u64;

core/src/ui/access.rs

@@ -0,0 +1,152 @@
+//! I4 (RUST.md): an AccessKit tree built from iris's own widget tree,
+//! shared by both backends -- `android/view.rs` pushes its `TreeUpdate`s
+//! through `accesskit_android::Adapter`, `default/mod.rs` through
+//! `accesskit_winit::Adapter`. Kept modular the way input's sense registry
+//! is: `Widgets::named()` is a side set populated only by `.label()`, so a
+//! widget nobody named is never visited here at all, not even to decide it
+//! has no name.
+//!
+//! The tree itself is deliberately flat -- one synthetic `Role::Window`
+//! root with every named widget as a direct child, in no particular order.
+//! iris's actual widget nesting (a label three `Span`s deep inside a
+//! `ScrollArea`) carries no accessibility meaning of its own here: nothing
+//! upstream of a named leaf needs a node, since a screen reader's own
+//! traversal (and uiautomator's tap-by-name, the pass condition this was
+//! built for) works from each node's on-screen bounds rather than from
+//! tree structure. Mirroring the real widget tree exactly would also mean
+//! rebuilding intermediate nodes whenever *any* container above a named
+//! widget resizes, which is most frames -- the flat shape is what keeps
+//! rebuilds tied to "a name, a role or a position actually changed".
+
+use crate::{PixelRegion, UiRenderState, UiRsc, WidgetId, Widgets, util::HashMap};
+use accesskit::{Node, NodeId, Rect, Role, TreeId, TreeInfo, TreeUpdate};
+
+/// Reserved for the synthetic root; every real widget's `SlotId::as_u64`
+/// starts at 1, so this can never collide with one (see that method's
+/// doc comment).
+const WINDOW_NODE: NodeId = NodeId(0);
+
+fn node_id(id: WidgetId) -> NodeId {
+    NodeId(id.as_u64())
+}
+
+#[derive(Clone, PartialEq)]
+struct Entry {
+    name: String,
+    role: Role,
+    bounds: PixelRegion,
+}
+
+fn entry_node(entry: &Entry) -> Node {
+    let mut node = Node::new(entry.role);
+    node.set_label(entry.name.clone());
+    node.set_bounds(Rect {
+        x0: entry.bounds.top_left.x as f64,
+        y0: entry.bounds.top_left.y as f64,
+        x1: entry.bounds.bot_right.x as f64,
+        y1: entry.bounds.bot_right.y as f64,
+    });
+    node
+}
+
+/// Owns the last tree pushed out, so `update` can tell "nothing
+/// accessibility-relevant changed" from "something did" without asking
+/// the platform adapter to diff two `Node`s itself. One of these per
+/// window/view -- `default::DefaultUiState` and `android::AndroidUiState`
+/// each keep one.
+#[derive(Default)]
+pub struct AccessTree {
+    known: HashMap<WidgetId, Entry>,
+    /// `TreeUpdate`s actually produced since the last `take_rebuilds` --
+    /// the AccessKit-tree twin of `UiRenderState::take_counters`. Should
+    /// stay at 0 across an unchanged frame and move by exactly 1 when a
+    /// named widget's position, name or role changes, however many other
+    /// widgets are on screen; see `iris/src/access_tests.rs`.
+    rebuilds: u64,
+}
+
+impl AccessTree {
+    pub fn new() -> Self {
+        Self::default()
+    }
+
+    fn collect(
+        widgets: &Widgets,
+        render: &UiRenderState,
+        rsc: &dyn UiRsc,
+    ) -> HashMap<WidgetId, Entry> {
+        let mut current = HashMap::default();
+        for id in widgets.named() {
+            let Some(bounds) = render.window_region(&id, rsc) else {
+                continue;
+            };
+            let Some(widget) = widgets.get_dyn(id) else {
+                continue;
+            };
+            current.insert(
+                id,
+                Entry {
+                    name: widgets.label(id).clone(),
+                    role: widget.access_role(),
+                    bounds,
+                },
+            );
+        }
+        current
+    }
+
+    /// Walks `widgets.named()`, looks up each one's current screen bounds
+    /// via `render.window_region` (which resolves the same move-chain
+    /// `resolved_region` does, so a moved subtree reports where it
+    /// actually is), and returns a full `TreeUpdate` if and only if that
+    /// set differs from the last call -- added, removed, renamed, or
+    /// moved/resized. A widget that is named but not currently active
+    /// (not drawn this frame) is left out, the same as one never named at
+    /// all.
+    pub fn update(
+        &mut self,
+        widgets: &Widgets,
+        render: &UiRenderState,
+        rsc: &dyn UiRsc,
+    ) -> Option<TreeUpdate> {
+        let current = Self::collect(widgets, render, rsc);
+        if current == self.known {
+            return None;
+        }
+        self.known = current.clone();
+        self.rebuilds += 1;
+        Some(build_update(&current))
+    }
+
+    /// The unconditional twin of `update`, for a platform adapter's
+    /// activation handler (`android/access.rs`'s `AndroidAccessSource`) --
+    /// AccessKit asks for a full tree the first time a client attaches,
+    /// which is exactly the case `update`'s diff-against-`known` is not
+    /// meant to answer (it may have already sent this same snapshot to a
+    /// client that has since detached and reattached).
+    pub fn build_full(widgets: &Widgets, render: &UiRenderState, rsc: &dyn UiRsc) -> TreeUpdate {
+        build_update(&Self::collect(widgets, render, rsc))
+    }
+
+    /// Reads and zeroes the rebuild counter, the same call shape as
+    /// `UiRenderState::take_counters`.
+    pub fn take_rebuilds(&mut self) -> u64 {
+        std::mem::take(&mut self.rebuilds)
+    }
+}
+
+fn build_update(current: &HashMap<WidgetId, Entry>) -> TreeUpdate {
+    let mut window = Node::new(Role::Window);
+    let mut nodes = Vec::with_capacity(current.len() + 1);
+    for (&id, entry) in current {
+        window.push_child(node_id(id));
+        nodes.push((node_id(id), entry_node(entry)));
+    }
+    nodes.push((WINDOW_NODE, window));
+    TreeUpdate {
+        nodes,
+        tree: Some(TreeInfo::new(WINDOW_NODE)),
+        tree_id: TreeId::ROOT,
+        focus: WINDOW_NODE,
+    }
+}

core/src/ui/active.rs

@@ -1,4 +1,6 @@
-use crate::{LayerId, MaskIdx, PrimitiveHandle, TextureHandle, UiRegion, WidgetId};
+use crate::{
+    LayerId, MaskIdx, MoveIdx, PrimitiveHandle, Size, TextureHandle, UiRegion, WidgetId, util::Vec2,
+};
 
 /// important non rendering data for retained drawing
 #[derive(Debug)]
@@ -9,6 +11,59 @@ pub struct ActiveData {
     pub textures: Vec<TextureHandle>,
     pub primitives: Vec<PrimitiveHandle>,
     pub children: Vec<WidgetId>,
+    /// The mask this widget was drawn **under** (its parent's), not the
+    /// one it set for itself -- see `own_mask` for that.
     pub mask: MaskIdx,
+    /// The mask slot this widget allocated for *itself* with
+    /// `Painter::set_mask`, or `MaskIdx::NONE`. Kept across redraws and
+    /// rewritten in place, the way `move_slot` is: a `Masked` that pushed
+    /// a fresh slot each draw left every already-drawn descendant --
+    /// which `draw_inner`'s unchanged-region fast path does not revisit --
+    /// clipping to the *old* slot's region, so a composer whose bar had
+    /// since been placed at the bottom of the screen was still being
+    /// clipped to a box at the top of it and drew nothing (measured
+    /// 2026-09-06: four mask entries live, none of them the widget's
+    /// current region). Its path out is the `undraw` branch of
+    /// `UiRenderState::remove`, which drops the self-ownership ref taken
+    /// when the slot was allocated.
+    pub own_mask: MaskIdx,
     pub layer: LayerId,
+    /// What `Widget::draw` returned the last time this widget was actually
+    /// drawn -- read by a parent placing this widget again without
+    /// redrawing it, replacing `Cache.size`'s old role. See LAYOUT.md
+    /// section 5.
+    pub size: Size,
+    /// This widget's slot in `UiData::move_offsets`, assigned on its first
+    /// draw and kept for the rest of its life (redraws reuse it in place
+    /// so a retained child's `parent` link never goes stale). See
+    /// LAYOUT.md section 2.
+    pub move_slot: MoveIdx,
+    /// How much of this widget's own `move_slot` delta is already folded
+    /// into `region` above, in window pixels. The two mechanisms that
+    /// write that slot disagree about this and cannot be told apart from
+    /// the slot alone: `UiRenderState::mov` shifts `region` and the delta
+    /// together (the *offered* region genuinely moved), while
+    /// `Painter::reposition` writes only the delta (`region` stays the
+    /// offered box and the delta says where inside it the content was
+    /// placed). So anything that wants the widget's real position --
+    /// `resolved_region`, and through it every hit test -- must subtract
+    /// this from the chain sum. Without it a panned widget's own hit box
+    /// sits at twice the pan while its descendants' are correct, which is
+    /// how it went unnoticed: the composer's field became untappable
+    /// after a finger pan (2026-09-06). Reset to zero whenever the widget
+    /// is really redrawn, since `draw_inner` zeroes the slot then too.
+    pub move_applied: Vec2,
+    /// The offset the last `Painter::reposition` placed this widget's
+    /// content at *within* `region`, in window pixels. The move slot has
+    /// exactly one owner and one meaning:
+    /// `move_offsets[move_slot] == move_applied + repositioned`. `mov`
+    /// adds to the first, `reposition` overwrites the second (it
+    /// recomputes `from` afresh every call, so repeating it must land on
+    /// the same answer rather than drifting), and both then rewrite the
+    /// slot from the sum -- which is what lets a parent both move a child
+    /// with its own layout and place it inside that moved region in one
+    /// frame. `LazySpan::place`'s Bottom-known branch does exactly that once a
+    /// row's blocks wrap. Reset to zero on a real redraw, with
+    /// `move_applied` and the slot itself.
+    pub repositioned: Vec2,
 }

core/src/ui/cache.rs

@@ -1,18 +0,0 @@
-use crate::{BothAxis, Len, UiVec2, WidgetId, util::HashMap};
-
-#[derive(Default)]
-pub struct Cache {
-    pub size: BothAxis<HashMap<WidgetId, (UiVec2, Len)>>,
-}
-
-impl Cache {
-    pub fn remove(&mut self, id: WidgetId) {
-        self.size.x.remove(&id);
-        self.size.y.remove(&id);
-    }
-
-    pub fn clear(&mut self) {
-        self.size.x.clear();
-        self.size.y.clear();
-    }
-}

core/src/ui/mod.rs

@@ -1,15 +1,16 @@
-use crate::{Mask, TextData, Textures, WeakWidget, WidgetId, Widgets, util::TrackedArena};
+use crate::{
+    Mask, MoveOffset, TextData, Textures, WeakWidget, WidgetId, Widgets, util::TrackedArena,
+};
 
+mod access;
 mod active;
-mod cache;
 mod painter;
 mod render_state;
-mod size;
 
+pub use access::*;
 pub use active::*;
 pub use painter::Painter;
 pub use render_state::*;
-pub use size::*;
 
 #[derive(Default)]
 pub struct UiData {
@@ -17,6 +18,52 @@ pub struct UiData {
     pub textures: Textures,
     pub text: TextData,
     pub masks: TrackedArena<Mask, u32>,
+    /// One entry per widget ever drawn, forming the parent-linked chain
+    /// `resolve_move` walks in both shader stages. Allocated once on a
+    /// widget's first draw and reused for every later redraw of the same
+    /// id (never reallocated), so a retained descendant's `parent` index
+    /// never goes stale -- see LAYOUT.md section 2.
+    pub move_offsets: TrackedArena<MoveOffset, u32>,
+    /// Every widget whose [`crate::Widget::tick`] should run before the
+    /// next frame -- today, a `LazySpan` coasting through a fling. Added by
+    /// [`Self::animate`] when the animation starts and removed by
+    /// [`Self::tick_animations`] the frame its `tick` answers `false`, so
+    /// a stopped animation costs nothing and a dropped widget cannot be
+    /// ticked (`get_dyn_mut` answers `None` and it is dropped the same
+    /// way).
+    animating: Vec<WidgetId>,
+}
+
+impl UiData {
+    /// Ask for `id`'s [`crate::Widget::tick`] to run every frame until it
+    /// says it is done. Idempotent -- registering an already-animating
+    /// widget is the ordinary case (a second fling before the first
+    /// settled) and must not tick it twice per frame.
+    pub fn animate(&mut self, id: WidgetId) {
+        if !self.animating.contains(&id) {
+            self.animating.push(id);
+        }
+    }
+
+    /// Tick every registered widget to `now`, drop the ones that finished,
+    /// and say whether any is still going -- which is a backend's cue to
+    /// ask for another frame. Called once per frame *before* the draw, so
+    /// what the frame draws is this instant's position rather than the
+    /// previous one's.
+    pub fn tick_animations(&mut self, now: std::time::Instant) -> bool {
+        // Taken out and put back rather than iterated in place: `tick`
+        // needs `&mut` on the widget arena this list lives beside, and a
+        // widget is free to register another one while ticking.
+        let mut registered = std::mem::take(&mut self.animating);
+        registered.retain(|&id| match self.widgets.get_dyn_mut(id) {
+            Some(widget) => widget.tick(now),
+            None => false,
+        });
+        for id in registered {
+            self.animate(id);
+        }
+        !self.animating.is_empty()
+    }
 }
 
 pub trait UiRsc {

core/src/ui/painter.rs

@@ -1,7 +1,10 @@
 use crate::{
-    Axis, Len, RenderedText, Size, SizeCtx, StrongWidget, TextAttrs, TextBuffer, TextData,
-    TextureHandle, UiRegion, UiRenderState, UiRsc, Widget, WidgetId,
-    render::{Mask, MaskIdx, 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,
 };
 
@@ -12,6 +15,11 @@ pub struct Painter<'a> {
 
     pub(super) region: UiRegion,
     pub(super) mask: MaskIdx,
+    pub(super) move_slot: MoveIdx,
+    /// This widget's own mask slot, reused across redraws -- see
+    /// `ActiveData::own_mask`. `MaskIdx::NONE` until `set_mask` is called
+    /// for the first time in this widget's life.
+    pub(super) own_mask: MaskIdx,
     pub(super) textures: Vec<TextureHandle>,
     pub(super) primitives: Vec<PrimitiveHandle>,
     pub(super) children: Vec<WidgetId>,
@@ -21,19 +29,49 @@ pub struct Painter<'a> {
 
 impl<'a> Painter<'a> {
     fn primitive_at<P: Primitive>(&mut self, primitive: P, region: UiRegion) {
-        let h = self.state.layers.write(
+        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,
                 region,
                 mask_idx: self.mask,
+                move_idx: self.move_slot,
             },
         );
         if self.mask != MaskIdx::NONE {
             // 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.own(h);
+        slot
+    }
+
+    /// Take ownership of a handle this widget just wrote.
+    ///
+    /// The one place a `PrimitiveHandle` enters `self.primitives`, and so
+    /// the one place that can keep `Primitives::handle_index` in step with
+    /// where it lands -- which is what `UiRenderState::apply_free` reads
+    /// instead of scanning this vec. Anything that writes a primitive
+    /// without coming through here leaves that index unset, and its
+    /// position in a layer's draw order stops being renumbered.
+    fn own(&mut self, h: PrimitiveHandle) {
+        self.state
+            .primitives
+            .set_handle_index(h.slot, self.primitives.len() as u32);
         self.primitives.push(h);
     }
 
@@ -46,75 +84,291 @@ impl<'a> Painter<'a> {
         self.primitive_at(primitive, region.within(&self.region));
     }
 
+    /// 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 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) {
-        assert!(self.mask == MaskIdx::NONE);
-        self.mask = self.rsc.ui_mut().masks.push(Mask { region });
+        let shape = self.write_primitive(RectPrimitive::color(Color::NONE), region, Drawn::No);
+        self.set_mask_to(shape);
     }
 
-    /// Draws a widget within this widget's region.
-    pub fn widget<W: ?Sized>(&mut self, id: &StrongWidget<W>) {
-        self.widget_at(id, self.region);
+    /// 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.
+        // Every build anybody runs here is release
+        // (docs/REVIEW-2026-09-07.md's R1).
+        assert!(
+            self.own_mask == MaskIdx::NONE || self.mask != self.own_mask,
+            "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 {
+            primitive: shape,
+            parent,
+        };
+        let old_parent = if self.own_mask == MaskIdx::NONE {
+            let slot = self.rsc.ui_mut().masks.push(mask);
+            // The one ref this widget holds on its own slot, so the slot
+            // outlives any single frame's primitives; released in
+            // `UiRenderState::remove`'s `undraw` branch.
+            self.rsc.ui_mut().masks.push_ref(slot);
+            self.own_mask = slot;
+            MaskIdx::NONE
+        } else {
+            let old = self.rsc.ui().masks[self.own_mask.idx()].parent;
+            *self.rsc.ui_mut().masks.get_mut(self.own_mask) = mask;
+            old
+        };
+        // The chain link's own ref, taken before the old one is dropped so
+        // that re-chaining to the same slot cannot free it in between.
+        // Released here when the link changes, and in
+        // `UiRenderState::remove` when this widget's slot goes.
+        if old_parent != parent {
+            if parent != MaskIdx::NONE {
+                self.rsc.ui_mut().masks.push_ref(parent);
+            }
+            if old_parent != MaskIdx::NONE {
+                self.rsc.ui_mut().masks.remove(old_parent);
+            }
+        }
+        self.mask = self.own_mask;
+    }
+
+    /// Draws a widget within this widget's region, returning the size it
+    /// reported using.
+    pub fn widget<W: ?Sized>(&mut self, id: &StrongWidget<W>) -> Size {
+        self.widget_at(id, self.region)
     }
 
     /// Draws a widget somewhere within this one.
     /// Useful for drawing child widgets in select areas.
-    pub fn widget_within<W: ?Sized>(&mut self, id: &StrongWidget<W>, region: UiRegion) {
-        self.widget_at(id, region.within(&self.region));
+    pub fn widget_within<W: ?Sized>(&mut self, id: &StrongWidget<W>, region: UiRegion) -> Size {
+        self.widget_at(id, region.within(&self.region))
     }
 
-    fn widget_at<W: ?Sized>(&mut self, id: &StrongWidget<W>, region: UiRegion) {
+    fn widget_at<W: ?Sized>(&mut self, id: &StrongWidget<W>, region: UiRegion) -> Size {
         self.children.push(id.id());
+        // Passed directly rather than looked up from `self.active`: this
+        // widget's own `ActiveData` (which would carry its `move_slot`) is
+        // not inserted there until *after* its own `Widget::draw` returns,
+        // so a lookup here -- for a child drawn partway through that same
+        // call -- would always find nothing. `self.move_slot` is this
+        // widget's own slot, already known, and always correct regardless
+        // of insertion order. See `UiRenderState::move_parent_of`.
         self.state.draw_inner(
             self.layer,
             id.id(),
             region,
             Some(self.id),
+            self.move_slot.idx() as u32,
             self.mask,
             None,
+            None,
+            crate::render::MaskIdx::NONE,
             self.rsc,
         );
+        self.state
+            .active
+            .get(&id.id())
+            .map(|a| a.size)
+            .unwrap_or_default()
+    }
+
+    /// Move an already-drawn child from wherever it currently sits to
+    /// `region` (resolved against this widget's own region, matching
+    /// `widget_within`) without a second draw -- an O(1) offset write via
+    /// `UiRenderState::mov`. For a container that draws a child
+    /// provisionally to learn its size (e.g. `Aligned`) and then places it
+    /// for real. Only valid when the target keeps the child's drawn size;
+    /// if the shape actually changes, the normal `widget_within` dispatch
+    /// (which detects that from the stored region) does the right thing
+    /// instead.
+    pub fn reposition<W: ?Sized>(&mut self, id: &StrongWidget<W>, region: UiRegion) {
+        let region = region.within(&self.region);
+        self.state.reposition(id.id(), region, self.rsc);
+    }
+
+    /// Draw `child` at a provisional region to learn its size under one
+    /// axis's worth of assumption, discard everything it wrote, then draw
+    /// it again at the region that assumption produced. For the rare
+    /// parent that cannot pick an offered size without already knowing the
+    /// answer. Twice the cost of one `draw`; every other case in this file
+    /// avoids it.
+    pub fn draw_twice<W: ?Sized>(
+        &mut self,
+        id: &StrongWidget<W>,
+        first: UiRegion,
+        second: impl FnOnce(Size) -> UiRegion,
+    ) -> Size {
+        let used = self.widget_within(id, first);
+        let region = second(used);
+        self.widget_within(id, region)
     }
 
     pub fn texture_within(&mut self, handle: &TextureHandle, region: UiRegion) {
         self.textures.push(handle.clone());
-        self.primitive_at(handle.primitive(), region.within(&self.region));
+        self.write_image(handle.image_index(), region.within(&self.region));
     }
 
     pub fn texture(&mut self, handle: &TextureHandle) {
         self.textures.push(handle.clone());
-        self.primitive(handle.primitive());
+        self.write_image(handle.image_index(), self.region);
     }
 
     pub fn texture_at(&mut self, handle: &TextureHandle, region: UiRegion) {
         self.textures.push(handle.clone());
-        self.primitive_at(handle.primitive(), region);
+        self.write_image(handle.image_index(), region);
     }
 
-    /// returns (handle, offset from top left)
-    pub fn render_text(&mut self, buffer: &mut TextBuffer, attrs: &TextAttrs) -> RenderedText {
+    /// A standalone image draws with its own bind group rather than sharing
+    /// the layer's one instanced draw, so it goes through
+    /// `Primitives::write_image` instead of `primitive_at`/`Primitive::vec`.
+    fn write_image(&mut self, texture_idx: u32, region: UiRegion) {
+        let h = self.state.write_image(
+            self.layer,
+            self.id,
+            texture_idx,
+            region,
+            self.mask,
+            self.move_slot,
+        );
+        if self.mask != MaskIdx::NONE {
+            self.rsc.ui_mut().masks.push_ref(self.mask);
+        }
+        self.own(h);
+    }
+
+    pub fn render_text(
+        &mut self,
+        buffer: &mut TextBuffer,
+        attrs: &TextAttrs,
+        width: Option<f32>,
+    ) -> RenderedText {
+        let density = self.state.density;
+        // Counted here rather than in `TextView::render`, which returns
+        // its memoized layout without reaching this -- so this counts
+        // shapes, not requests. `UiRenderState::take_counters`.
+        self.state.shape_count += 1;
         let ui = self.rsc.ui_mut();
-        ui.text.draw(buffer, attrs, &mut ui.textures)
+        ui.text
+            .render(buffer, attrs, width, &mut ui.textures, density)
+    }
+
+    /// Which glyph atlas the glyphs handed out right now belong to --
+    /// what a widget caching a [`RenderedText`] across frames has to
+    /// compare against before re-emitting it (`GlyphAtlas::clear`).
+    pub fn atlas_generation(&mut self) -> u64 {
+        self.rsc.ui_mut().text.atlas.generation()
+    }
+
+    /// Draw a laid-out string: one quad per glyph, all sampling the atlas.
+    ///
+    /// `origin` is where the text's top-left goes; every glyph is placed at an
+    /// absolute pixel offset from it, so re-drawing after a resize is this loop
+    /// and nothing else.
+    pub fn glyphs(&mut self, text: &RenderedText, origin: UiRegion) {
+        // A caller re-emitting quads placed against an atlas that has since
+        // been cleared draws every glyph from coordinates now holding
+        // something else. Caught at the submission rather than on screen,
+        // where it reads as fragments of unrelated letters. `assert_eq!`
+        // for R1's reason: two integers per laid-out string, not per
+        // glyph, and the failure is unreadable text on a release build.
+        assert_eq!(
+            text.generation,
+            self.atlas_generation(),
+            "glyphs placed against atlas generation {} submitted against {}: the holder did not \
+             re-render after the atlas was cleared",
+            text.generation,
+            self.atlas_generation(),
+        );
+        let flags_for = |is_color| {
+            if is_color {
+                GlyphPrimitive::IS_COLOR
+            } else {
+                0
+            }
+        };
+        for glyph in text.glyphs.iter() {
+            let mut region = origin;
+            region.x.end = region.x.start;
+            region.y.end = region.y.start;
+            let mut region = region.offset(UiVec2::abs(glyph.offset));
+            region.x.end = region.x.start + UiScalar::abs(glyph.entry.width as f32);
+            region.y.end = region.y.start + UiScalar::abs(glyph.entry.height as f32);
+            self.primitive_at(
+                GlyphPrimitive::new(
+                    glyph.entry.uv_min,
+                    glyph.entry.uv_max,
+                    glyph.entry.layer,
+                    glyph.color,
+                    flags_for(glyph.entry.is_color),
+                ),
+                region,
+            );
+        }
     }
 
     pub fn region(&self) -> UiRegion {
         self.region
     }
 
-    pub fn size<W: ?Sized + Widget>(&mut self, id: &StrongWidget<W>) -> Size {
-        self.size_ctx().size(id)
-    }
-
-    pub fn len_axis<W: ?Sized + Widget>(&mut self, id: &StrongWidget<W>, axis: Axis) -> Len {
-        match axis {
-            Axis::X => self.size_ctx().width(id),
-            Axis::Y => self.size_ctx().height(id),
-        }
-    }
-
     pub fn output_size(&self) -> Vec2 {
         self.state.output_size
     }
 
+    /// Physical pixels per `dp` -- see `UiRenderState::density`'s field
+    /// doc. What `Len::dp`'s `apply_rest` call resolves against.
+    pub fn density(&self) -> f32 {
+        self.state.density
+    }
+
     pub fn px_size(&mut self) -> Vec2 {
         self.region.size().to_abs(self.state.output_size)
     }
@@ -138,8 +392,4 @@ impl<'a> Painter<'a> {
     pub fn id(&self) -> &WidgetId {
         &self.id
     }
-
-    pub fn size_ctx(&mut self) -> SizeCtx<'_> {
-        self.state.size_ctx(self.id, self.region.size(), self.rsc)
-    }
 }

core/src/ui/size.rs

@@ -1,86 +0,0 @@
-use crate::{
-    Axis, AxisT, IdLike, Len, RenderedText, Size, TextAttrs, TextBuffer, TextData, Textures,
-    UiVec2, WidgetAxisFns, WidgetId, Widgets, XAxis, YAxis, ui::cache::Cache, util::Vec2,
-};
-
-pub struct SizeCtx<'a> {
-    pub text: &'a mut TextData,
-    pub textures: &'a mut Textures,
-    pub(super) source: WidgetId,
-    pub(super) widgets: &'a Widgets,
-    pub(super) cache: &'a mut Cache,
-    /// TODO: should this be pub? rn used for sized
-    pub outer: UiVec2,
-    pub(super) output_size: Vec2,
-    pub(super) id: WidgetId,
-}
-
-impl SizeCtx<'_> {
-    pub fn id(&self) -> &WidgetId {
-        &self.id
-    }
-
-    pub fn source(&self) -> &WidgetId {
-        &self.source
-    }
-
-    pub(super) fn len_inner<A: const AxisT>(&mut self, id: WidgetId) -> Len {
-        if let Some((_, len)) = self.cache.size.axis::<A>().get(&id) {
-            return *len;
-        }
-        let len = self
-            .widgets
-            .get_dyn_dynamic(id)
-            .desired_len::<A>(&mut SizeCtx {
-                text: self.text,
-                textures: self.textures,
-                source: self.source,
-                widgets: self.widgets,
-                cache: self.cache,
-                outer: self.outer,
-                output_size: self.output_size,
-                id,
-            });
-        self.cache.size.axis::<A>().insert(id, (self.outer, len));
-        len
-    }
-
-    pub fn width(&mut self, id: impl IdLike) -> Len {
-        self.len_inner::<XAxis>(id.id())
-    }
-
-    pub fn height(&mut self, id: impl IdLike) -> Len {
-        self.len_inner::<YAxis>(id.id())
-    }
-
-    pub fn len_axis(&mut self, id: impl IdLike, axis: Axis) -> Len {
-        match axis {
-            Axis::X => self.width(id),
-            Axis::Y => self.height(id),
-        }
-    }
-
-    pub fn size(&mut self, id: impl IdLike) -> Size {
-        let id = id.id();
-        Size {
-            x: self.width(id),
-            y: self.height(id),
-        }
-    }
-
-    pub fn px_size(&mut self) -> Vec2 {
-        self.outer.to_abs(self.output_size)
-    }
-
-    pub fn output_size(&mut self) -> Vec2 {
-        self.output_size
-    }
-
-    pub fn draw_text(&mut self, buffer: &mut TextBuffer, attrs: &TextAttrs) -> RenderedText {
-        self.text.draw(buffer, attrs, self.textures)
-    }
-
-    pub fn label(&self, id: WidgetId) -> &String {
-        self.widgets.label(id)
-    }
-}

core/src/util/arena.rs

@@ -71,6 +71,15 @@ impl<T, I: IdNum> TrackedArena<T, I> {
         self.refs[i.idx()] += 1;
     }
 
+    /// Mutable access to an existing entry, for the rare case (the move
+    /// offset chain) where an already-allocated slot is updated in place
+    /// rather than replaced. Marks the arena changed so the GPU copy is
+    /// re-uploaded.
+    pub fn get_mut(&mut self, id: Id<I>) -> &mut T {
+        self.changed = true;
+        &mut self.inner.data[id.idx()]
+    }
+
     pub fn remove(&mut self, id: Id<I>) -> T
     where
         T: Copy,

core/src/util/math.rs

@@ -9,15 +9,16 @@ pub const trait DivOr {
     fn div_or(self, rhs: Self, other: Self) -> Self;
 }
 
-impl const DivOr for f32 {
+const impl DivOr for f32 {
     fn div_or(self, rhs: Self, other: Self) -> Self {
         let res = self / rhs;
         if res.is_nan() { other } else { res }
     }
 }
 
-impl<T: const Add<Output = T> + const Sub<Output = T> + const Mul<Output = T> + const DivOr + Copy> const
-    LerpUtil for T
+const impl<
+    T: const Add<Output = T> + const Sub<Output = T> + const Mul<Output = T> + const DivOr + Copy,
+> LerpUtil for T
 {
     /// linear interpolation
     /// from * (1.0 - self) + to * self
@@ -37,7 +38,7 @@ macro_rules! impl_op {
             use super::*;
             #[allow(unused_imports)]
             use std::ops::*;
-            impl const $op for $T {
+            const impl $op for $T {
                 type Output = Self;
 
                 fn $fn(self, rhs: Self) -> Self::Output {
@@ -46,12 +47,12 @@ macro_rules! impl_op {
                     }
                 }
             }
-            impl const $opa for $T {
+            const impl $opa for $T {
                 fn $fna(&mut self, rhs: Self) {
                     *self = self.$fn(rhs);
                 }
             }
-            impl const $op<f32> for $T {
+            const impl $op<f32> for $T {
                 type Output = Self;
 
                 fn $fn(self, rhs: f32) -> Self::Output {
@@ -60,7 +61,7 @@ macro_rules! impl_op {
                     }
                 }
             }
-            impl const $op<$T> for f32 {
+            const impl $op<$T> for f32 {
                 type Output = $T;
 
                 fn $fn(self, rhs: $T) -> Self::Output {
@@ -69,7 +70,7 @@ macro_rules! impl_op {
                     }
                 }
             }
-            impl const $opa<f32> for $T {
+            const impl $opa<f32> for $T {
                 fn $fna(&mut self, rhs: f32) {
                     *self = self.$fn(rhs);
                 }

core/src/util/slot.rs

@@ -4,6 +4,17 @@ pub struct SlotId {
     genr: u32,
 }
 
+impl SlotId {
+    /// A stable, collision-free `u64` encoding of this id -- for a caller
+    /// (accesskit's `NodeId`, today) that wants a flat integer key rather
+    /// than the two `u32`s. `idx` is offset by one so no real id ever
+    /// encodes to 0, which callers can then reserve for their own
+    /// out-of-band root/window node.
+    pub fn as_u64(&self) -> u64 {
+        ((self.idx as u64) + 1) << 32 | self.genr as u64
+    }
+}
+
 pub struct SlotVec<T> {
     data: Vec<(u32, Option<T>)>,
     free: Vec<u32>,

core/src/util/vec2.rs

@@ -67,7 +67,7 @@ impl_op!(Vec2 Sub sub; x y);
 impl_op!(Vec2 Mul mul; x y);
 impl_op!(Vec2 Div div; x y);
 
-impl const DivOr for Vec2 {
+const impl DivOr for Vec2 {
     fn div_or(self, rhs: Self, other: Self) -> Self {
         Self {
             x: self.x.div_or(rhs.x, other.x),

core/src/widget/mod.rs

@@ -1,4 +1,4 @@
-use crate::{Axis, AxisT, Len, Painter, SizeCtx};
+use crate::{Painter, Size};
 use std::any::Any;
 
 mod data;
@@ -16,31 +16,59 @@ pub use view::*;
 pub use widgets::*;
 
 pub trait Widget: Any {
-    fn draw(&mut self, painter: &mut Painter);
-    fn desired_width(&mut self, ctx: &mut SizeCtx) -> Len;
-    fn desired_height(&mut self, ctx: &mut SizeCtx) -> Len;
-}
+    /// Draw within `painter.region()` (the space the parent offered) and
+    /// report how much of it was actually used, per axis.
+    fn draw(&mut self, painter: &mut Painter) -> Size;
 
-pub trait WidgetAxisFns {
-    fn desired_len<A: AxisT>(&mut self, ctx: &mut SizeCtx) -> Len;
-}
+    /// True if `draw`'s output (both the primitives it writes and the
+    /// `Size` it returns) is the same for any `painter.region()` of the
+    /// same *content* -- an icon, a fixed-size rect, an already-decoded
+    /// image at its natural size. Default `false` (redraw on any change to
+    /// the offered region) because assuming independence wrongly produces
+    /// a stale draw; a widget must opt in. See LAYOUT.md.
+    fn is_size_independent(&self) -> bool {
+        false
+    }
 
-impl<W: Widget + ?Sized> WidgetAxisFns for W {
-    fn desired_len<A: AxisT>(&mut self, ctx: &mut SizeCtx) -> Len {
-        match A::get() {
-            Axis::X => self.desired_width(ctx),
-            Axis::Y => self.desired_height(ctx),
-        }
+    /// What kind of control this is, for the AccessKit tree `ui::access`
+    /// builds (RUST.md's I4). Only consulted for a widget that also has an
+    /// explicit `.label()` -- an unnamed widget is never visited by that
+    /// tree at all, named or not, so the default here costs nothing except
+    /// at the handful of call sites that opt in. Default `Unknown` (a
+    /// generic control with no more specific semantics); a widget with a
+    /// real platform equivalent -- `TextEdit`'s `MultilineTextInput` --
+    /// overrides it.
+    fn access_role(&self) -> accesskit::Role {
+        accesskit::Role::Unknown
+    }
+
+    /// Advance whatever this widget is animating to `now`, and say whether
+    /// it is still animating afterwards. Default: nothing is, so a widget
+    /// opts in by overriding this *and* by something calling
+    /// [`crate::UiData::animate`] with its id when the animation starts --
+    /// which is that animation's path out, since the driver
+    /// ([`crate::UiData::tick_animations`]) drops every id whose `tick`
+    /// answers `false`.
+    ///
+    /// Called once per frame, before the frame's draw, by whichever
+    /// backend owns the surface; a `true` answer is what makes that
+    /// backend ask for another frame. So this is the only thing in iris
+    /// that moves without an input event, and a widget that animates
+    /// without registering simply never moves -- which is exactly how a
+    /// finger fling looked on Iris's phone before this existed.
+    #[allow(unused_variables)]
+    fn tick(&mut self, now: std::time::Instant) -> bool {
+        false
     }
 }
 
 impl Widget for () {
-    fn draw(&mut self, _: &mut Painter) {}
-    fn desired_width(&mut self, _: &mut SizeCtx) -> Len {
-        Len::ZERO
+    fn draw(&mut self, _: &mut Painter) -> Size {
+        Size::ZERO
     }
-    fn desired_height(&mut self, _: &mut SizeCtx) -> Len {
-        Len::ZERO
+
+    fn is_size_independent(&self) -> bool {
+        true
     }
 }
 

core/src/widget/widgets.rs

@@ -11,6 +11,11 @@ pub struct Widgets {
     send: Sender<WidgetId>,
     recv: Receiver<WidgetId>,
     pub(crate) waiting: HashSet<WidgetId>,
+    /// Every widget that has ever been given an explicit `.label()` --
+    /// `ui::access::AccessTree` walks exactly this set, not the whole
+    /// arena, so a widget nobody named costs it nothing. Symmetric with
+    /// `free_next` below, which is this set's one removal path.
+    named: HashSet<WidgetId>,
 }
 
 impl Widgets {
@@ -20,6 +25,7 @@ impl Widgets {
             needs_redraw: Default::default(),
             vec: Default::default(),
             waiting: Default::default(),
+            named: Default::default(),
             send,
             recv,
         }
@@ -95,9 +101,20 @@ impl Widgets {
         &self.data(id.id()).unwrap().label
     }
 
-    /// useful for debugging
+    /// Also the one place a widget opts into `ui::access`'s AccessKit tree
+    /// (RUST.md's I4) -- see `named`'s doc comment.
     pub fn set_label(&mut self, id: impl IdLike, label: String) {
-        self.data_mut(id.id()).unwrap().label = label;
+        let id = id.id();
+        self.data_mut(id).unwrap().label = label;
+        self.named.insert(id);
+    }
+
+    /// Every widget with an explicit name, for `ui::access::AccessTree` to
+    /// walk. Order is unspecified; `AccessTree` doesn't need one; a screen
+    /// reader's own traversal is worked out by uiautomator from each
+    /// node's on-screen bounds instead.
+    pub fn named(&self) -> impl Iterator<Item = WidgetId> + '_ {
+        self.named.iter().copied()
     }
 
     pub fn data_mut(&mut self, id: impl IdLike) -> Option<&mut WidgetData> {
@@ -107,6 +124,7 @@ impl Widgets {
     pub fn free_next(&mut self) -> Option<WidgetId> {
         let next = self.recv.try_recv().ok()?;
         self.vec.free(next);
+        self.named.remove(&next);
         Some(next)
     }