diff --git a/crates/gpui/Cargo.toml b/crates/gpui/Cargo.toml index 2e9350d64a..c46fce8fb9 100644 --- a/crates/gpui/Cargo.toml +++ b/crates/gpui/Cargo.toml @@ -215,6 +215,10 @@ path = "examples/learn/text.rs" name = "transition" path = "examples/learn/transition.rs" +[[example]] +name = "blur" +path = "examples/learn/blur.rs" + # ============================================================================ # Bench Examples - Performance benchmarks # ============================================================================ diff --git a/crates/gpui/examples/learn/blur.rs b/crates/gpui/examples/learn/blur.rs new file mode 100644 index 0000000000..bd4270bde8 --- /dev/null +++ b/crates/gpui/examples/learn/blur.rs @@ -0,0 +1,313 @@ +//! Blur Filters Example +//! +//! Demonstrates the two CSS-style blur filters: +//! +//! 1. `backdrop_blur` — frosted glass: blurs whatever is rendered behind the element. +//! Shown as a translucent panel, and again inside a `deferred()` popover (to prove the +//! backdrop snapshot includes everything beneath an overlay, and that the overlay sorts on top). +//! 2. `blur` — content blur: blurs the element and its own children as a group. +//! Shown with text and again with a row of colored chips. +//! +//! It also stresses two content-blur edge cases: +//! +//! 3. Nested content blur — a `blur()` element inside another `blur()` element, so the inner +//! subtree is blurred twice (its own filter, then again as part of the outer group). This +//! exercises the renderer's per-nesting-level group textures. +//! 4. Adjacent blocks with no gap on a dark parent, shown two ways: `blur` on each block (every +//! block is its own group, so the parent shows through each seam — exactly like CSS `filter` +//! on each sibling) versus `blur` on the parent (one group covering all blocks, so the blur is +//! continuous and the seams are clean — the CSS "blur the wrapper" idiom). + +use gpui::{ + App, Bounds, Context, Render, Window, WindowBounds, WindowOptions, deferred, div, point, + prelude::*, px, rgb, rgba, size, +}; + +struct BlurExample; + +/// A vivid, gap-free background so blur is obvious: a grid of saturated tiles filling the window. +fn busy_background() -> impl IntoElement { + let palette = [ + 0xef4444, 0xf97316, 0xeab308, 0x22c55e, 0x06b6d4, 0x3b82f6, 0x8b5cf6, 0xec4899, + ]; + let mut next = 0usize; + div() + .absolute() + .inset_0() + .flex() + .flex_col() + .children((0..9).map(|_| { + div().flex().flex_1().children( + (0..8) + .map(|_| { + let hex = palette[next % palette.len()]; + next += 1; + div() + .flex_1() + .h_full() + .bg(rgb(hex)) + .flex() + .items_center() + .justify_center() + .text_color(rgb(0xffffff)) + .text_xl() + .child("◆") + }) + .collect::>(), + ) + })) +} + +/// A translucent rounded panel that frosts the content behind it. +fn frosted_panel() -> impl IntoElement { + div() + .absolute() + .left(px(60.)) + .top(px(120.)) + .w(px(360.)) + .h(px(200.)) + .rounded_xl() + .bg(rgba(0xffffff30)) + .backdrop_blur(px(24.)) + .border_1() + .border_color(rgba(0xffffff60)) + .flex() + .items_center() + .justify_center() + .text_color(rgb(0x111111)) + .text_2xl() + .child("backdrop_blur(24px)") +} + +/// A popover painted via `deferred()` so it sits above everything; its backdrop blur must +/// still pick up the panel and background beneath it. +fn deferred_popover() -> impl IntoElement { + deferred( + div() + .absolute() + .left(px(260.)) + .top(px(260.)) + .w(px(300.)) + .h(px(150.)) + .rounded_lg() + .bg(rgba(0x1e293b66)) + .backdrop_blur(px(12.)) + .border_1() + .border_color(rgba(0xffffffaa)) + .flex() + .items_center() + .justify_center() + .text_color(rgb(0xffffff)) + .text_xl() + .child("deferred + backdrop_blur"), + ) +} + +/// A self-blurred element (CSS `filter: blur`) — its own content is blurred as a group. +fn content_blurred() -> impl IntoElement { + div() + .absolute() + .left(px(120.)) + .top(px(420.)) + .w(px(280.)) + .h(px(120.)) + .blur(px(5.)) + .bg(rgb(0x0f172a)) + .rounded_lg() + .flex() + .items_center() + .justify_center() + .text_color(rgb(0xfacc15)) + .text_3xl() + .child("blur(5px) content") +} + +/// A content-blurred element with richer content (a row of colored chips), so the `filter: blur` +/// effect — the element and its children blurred as one group — is clearly visible. +fn content_blurred_rich() -> impl IntoElement { + div() + .absolute() + .left(px(120.)) + .top(px(580.)) + .w(px(280.)) + .h(px(120.)) + .blur(px(6.)) + .bg(rgb(0x1e293b)) + .rounded_lg() + .flex() + .items_center() + .justify_center() + .gap_3() + .children([0xef4444, 0x22c55e, 0x3b82f6].into_iter().map(|hex| { + div().w(px(48.)).h(px(48.)).rounded_md().bg(rgb(hex)) + })) +} + +/// Nested content blur: a `blur()` element inside another `blur()` element. The inner block is +/// blurred by its own filter and then again as part of the outer group, exercising the renderer's +/// per-nesting-level isolated group textures (up to `MAX_FILTER_DEPTH`). The inner content should +/// read as markedly softer than the outer block's own text. +fn nested_content_blurred() -> impl IntoElement { + div() + .absolute() + .left(px(740.)) + .top(px(80.)) + .w(px(290.)) + .h(px(220.)) + .blur(px(3.)) + .bg(rgb(0x1e293b)) + .rounded_xl() + .flex() + .flex_col() + .items_center() + .justify_center() + .gap_4() + .text_color(rgb(0xe2e8f0)) + .text_xl() + .child("outer blur(3px)") + .child( + div() + .w(px(190.)) + .h(px(100.)) + .blur(px(8.)) + .bg(rgb(0xf59e0b)) + .rounded_lg() + .flex() + .items_center() + .justify_center() + .text_color(rgb(0x111111)) + .text_2xl() + .child("inner blur(8px)"), + ) +} + +const SEAM_COLORS: [u32; 4] = [0xef4444, 0x22c55e, 0x3b82f6, 0xeab308]; + +/// One numbered, brightly-coloured block of the seam row. With `blur_each` it becomes its own +/// content-filter group; otherwise it is a plain block (relying on a blurred parent, if any). +/// Square corners on purpose: gpui content masks are axis-aligned rectangles, so a *rounded* +/// parent would not clip the blurred children to its radius and the busy background would leak +/// through the corner triangles — a separate concern from the seam blending under test here. +fn seam_block(i: usize, hex: u32, blur_each: bool) -> impl IntoElement { + let block = div().flex_1().h_full(); + let block = if blur_each { block.blur(px(5.)) } else { block }; + block + .bg(rgb(hex)) + .flex() + .items_center() + .justify_center() + .text_color(rgb(0xffffff)) + .text_2xl() + .child(format!("{}", i + 1)) +} + +/// Adjacent blocks, each its OWN content-filter group (`blur` on every block), no gap, dark parent. +/// This is CSS `filter: blur()` on each sibling: every block fades to transparent at its edges and +/// composites independently, so the dark parent shows through each seam by roughly `α_left · α_right` +/// (peaking at ~25% right on the seam). This matches the web — the clean alternative is the next panel. +fn adjacent_per_block_blur() -> impl IntoElement { + div() + .absolute() + .left(px(740.)) + .top(px(350.)) + .w(px(290.)) + .h(px(110.)) + .bg(rgb(0x050505)) + .flex() + .children( + SEAM_COLORS + .into_iter() + .enumerate() + .map(|(i, hex)| seam_block(i, hex, true)), + ) +} + +/// The same adjacent blocks, but `blur` is on the PARENT — one content-filter group covering all +/// four. The blocks are opaque and touching, so the group's interior has no transparency: the blur +/// is continuous across the seams and only the group's outer edge fades. This is the CSS "blur the +/// wrapper, not each child" idiom, and the seams come out clean. +fn adjacent_group_blur() -> impl IntoElement { + div() + .absolute() + .left(px(740.)) + .top(px(510.)) + .w(px(290.)) + .h(px(110.)) + .bg(rgb(0x050505)) + .blur(px(5.)) + .flex() + .children( + SEAM_COLORS + .into_iter() + .enumerate() + .map(|(i, hex)| seam_block(i, hex, false)), + ) +} + +/// A small dark pill label so the two new test sections are identifiable over the busy background. +fn caption(text: &'static str, left: f32, top: f32) -> impl IntoElement { + div() + .absolute() + .left(px(left)) + .top(px(top)) + .px_2() + .py_1() + .rounded_md() + .bg(rgba(0x000000cc)) + .text_color(rgb(0xffffff)) + .text_sm() + .child(text) +} + +impl Render for BlurExample { + fn render(&mut self, _: &mut Window, _: &mut Context) -> impl IntoElement { + div() + .relative() + .size_full() + .bg(rgb(0x000000)) + .child(busy_background()) + .child(frosted_panel()) + .child(content_blurred()) + .child(content_blurred_rich()) + .child(nested_content_blurred()) + .child(adjacent_per_block_blur()) + .child(adjacent_group_blur()) + .child(caption("nested content blur", 740., 50.)) + .child(caption( + "adjacent — blur each block (seams, = CSS)", + 740., + 322., + )) + .child(caption( + "adjacent — blur the parent (one group, clean)", + 740., + 482., + )) + .child(deferred_popover()) + } +} + +fn main() { + gpui_platform::application().run(|cx: &mut App| { + cx.activate(true); + cx.on_window_closed(|cx, _| { + if cx.windows().is_empty() { + cx.quit(); + } + }) + .detach(); + + let bounds = Bounds { + origin: point(px(100.), px(100.)), + size: size(px(1060.), px(760.)), + }; + cx.open_window( + WindowOptions { + window_bounds: Some(WindowBounds::Windowed(bounds)), + ..Default::default() + }, + |_, cx| cx.new(|_| BlurExample), + ) + .expect("failed to open window"); + }); +} diff --git a/crates/gpui/src/bounds_tree.rs b/crates/gpui/src/bounds_tree.rs index e95e32477b..96bebda168 100644 --- a/crates/gpui/src/bounds_tree.rs +++ b/crates/gpui/src/bounds_tree.rs @@ -28,6 +28,10 @@ where root: Option, /// Index of the leaf with the highest ordering (for fast-path lookups). max_leaf: Option, + /// Minimum ordering assigned to any subsequent insert. Raised before painting deferred + /// draws so overlays always sort above the main scene (and their orders can't fall inside a + /// content-filter order range from the main scene). 0 means no floor. + order_floor: u32, /// Reusable stack for tree traversal during insertion. insert_path: Vec, /// Reusable stack for search operations. @@ -109,10 +113,34 @@ where self.nodes.clear(); self.root = None; self.max_leaf = None; + self.order_floor = 0; self.insert_path.clear(); self.search_stack.clear(); } + /// Raise the minimum ordering for subsequent inserts to `floor`. Relative ordering above the + /// floor is preserved (overlapping inserts still step above one another). + pub fn set_order_floor(&mut self, floor: u32) { + self.order_floor = self.order_floor.max(floor); + } + + /// The highest ordering assigned to any bounds so far (0 if empty). + pub fn max_order(&self) -> u32 { + self.max_leaf.map_or(0, |idx| self.nodes[idx].max_order) + } + + /// Inserts bounds with an ordering strictly greater than *every* existing bounds (not just + /// intersecting ones), and returns that ordering. Used for content-filter group boundaries + /// (which must sort after all previously-painted content so their order range can't collide + /// with unrelated non-overlapping content that reuses low orderings) and to raise the order + /// floor before painting deferred draws (so overlays always sort above the main scene). + pub fn insert_above_all(&mut self, new_bounds: Bounds) -> u32 { + let ordering = self.max_order() + 1; + let new_leaf_idx = self.insert_leaf(new_bounds, ordering); + self.max_leaf = Some(new_leaf_idx); + ordering + } + /// Inserts bounds into the tree and returns its assigned ordering. /// /// The ordering is one greater than the maximum ordering of any @@ -120,7 +148,7 @@ where pub fn insert(&mut self, new_bounds: Bounds) -> u32 { // Find maximum ordering among intersecting bounds let max_intersecting = self.find_max_ordering(&new_bounds); - let ordering = max_intersecting + 1; + let ordering = (max_intersecting + 1).max(self.order_floor); // Insert the new leaf let new_leaf_idx = self.insert_leaf(new_bounds, ordering); @@ -365,6 +393,7 @@ where nodes: Vec::new(), root: None, max_leaf: None, + order_floor: 0, insert_path: Vec::new(), search_stack: Vec::new(), } diff --git a/crates/gpui/src/scene.rs b/crates/gpui/src/scene.rs index e1e987f3df..93a4299e9a 100644 --- a/crates/gpui/src/scene.rs +++ b/crates/gpui/src/scene.rs @@ -6,8 +6,9 @@ use serde::{Deserialize, Serialize}; use crate::{ AtlasTextureId, AtlasTile, Background, Bounds, ContentMask, Corners, Edges, Hsla, Pixels, - Point, Radians, ScaledPixels, Size, bounds_tree::BoundsTree, point, + Point, Radians, ScaledFilter, ScaledPixels, Size, bounds_tree::BoundsTree, point, }; +use smallvec::SmallVec; use std::{ fmt::Debug, iter::Peekable, @@ -36,6 +37,8 @@ pub struct Scene { pub subpixel_sprites: Vec, pub polychrome_sprites: Vec, pub surfaces: Vec, + pub backdrop_filters: Vec, + pub filter_boundaries: Vec, } #[expect(missing_docs)] @@ -52,6 +55,8 @@ impl Scene { self.subpixel_sprites.clear(); self.polychrome_sprites.clear(); self.surfaces.clear(); + self.backdrop_filters.clear(); + self.filter_boundaries.clear(); } pub fn len(&self) -> usize { @@ -70,21 +75,50 @@ impl Scene { self.paint_operations.push(PaintOperation::EndLayer); } + /// Raise the draw-order floor so every primitive inserted afterwards sorts above everything + /// inserted before. Called before painting deferred draws so overlays (tooltips, popovers, + /// drag images) sort above the main scene — and a deferred backdrop's order can't fall inside + /// a content-filter (`filter`) order range left behind by the main scene. + pub fn raise_order_floor(&mut self) { + let floor = self.primitive_bounds.max_order() + 1; + self.primitive_bounds.set_order_floor(floor); + } + pub fn insert_primitive(&mut self, primitive: impl Into) { let mut primitive = primitive.into(); let clipped_bounds = primitive .bounds() .intersect(&primitive.content_mask().bounds); - if clipped_bounds.is_empty() { + // Content-filter boundaries must always be inserted as matched pairs — dropping one + // (e.g. for an empty clipped region) would orphan its partner and corrupt the renderer's + // target stack. Each marker takes an order strictly above ALL prior content, so the start + // sorts after everything painted before it and the element's own children (which overlap + // the marker bounds) sort strictly above the start. This keeps a marker's order range from + // colliding with unrelated non-overlapping content that reuses low orderings (e.g. a + // background grid), which would otherwise sweep that content into the group. Content + // painted *after* the group is held above it by raising the order floor when the end + // marker is inserted (see below) — otherwise a later non-overlapping sibling could reuse a + // low order that lands inside the start..end range and be swept into the group. + let is_filter_boundary = matches!(primitive, Primitive::FilterBoundary(_)); + + if clipped_bounds.is_empty() && !is_filter_boundary { return; } - let order = self - .layer_stack - .last() - .copied() - .unwrap_or_else(|| self.primitive_bounds.insert(clipped_bounds)); + let order = if is_filter_boundary { + let order_bounds = if clipped_bounds.is_empty() { + *primitive.bounds() + } else { + clipped_bounds + }; + self.primitive_bounds.insert_above_all(order_bounds) + } else { + self.layer_stack + .last() + .copied() + .unwrap_or_else(|| self.primitive_bounds.insert(clipped_bounds)) + }; match &mut primitive { Primitive::Shadow(shadow) => { shadow.order = order; @@ -119,6 +153,22 @@ impl Scene { surface.order = order; self.surfaces.push(surface.clone()); } + Primitive::BackdropFilter(filter) => { + filter.order = order; + self.backdrop_filters.push(filter.clone()); + } + Primitive::FilterBoundary(boundary) => { + boundary.order = order; + if !boundary.is_start { + // A closed content-filter group is a draw-order barrier: everything painted + // afterwards must sort above the group's end marker so it can't fall back + // inside the group's order range (subsequent non-overlapping content otherwise + // reuses a low order). Mirrors the floor raised before deferred draws in + // `raise_order_floor`. + self.primitive_bounds.set_order_floor(order + 1); + } + self.filter_boundaries.push(boundary.clone()); + } } self.paint_operations .push(PaintOperation::Primitive(primitive)); @@ -146,6 +196,13 @@ impl Scene { self.polychrome_sprites .sort_by_key(|sprite| (sprite.order, sprite.tile.tile_id)); self.surfaces.sort_by_key(|surface| surface.order); + self.backdrop_filters.sort_by_key(|filter| filter.order); + // Markers normally get distinct, monotonically-increasing orders (children overlap + // their group bounds and so sort strictly between the start and end). The `!is_start` + // tiebreak only matters for a degenerate empty group whose start and end tie: it keeps + // the start (false = 0) ahead of the end (true = 1) so the pair stays well-formed. + self.filter_boundaries + .sort_by_key(|boundary| (boundary.order, !boundary.is_start)); } #[cfg_attr( @@ -173,6 +230,10 @@ impl Scene { polychrome_sprites_iter: self.polychrome_sprites.iter().peekable(), surfaces_start: 0, surfaces_iter: self.surfaces.iter().peekable(), + backdrop_filters_start: 0, + backdrop_filters_iter: self.backdrop_filters.iter().peekable(), + filter_boundaries_start: 0, + filter_boundaries_iter: self.filter_boundaries.iter().peekable(), } } } @@ -186,6 +247,9 @@ impl Scene { allow(dead_code) )] pub(crate) enum PrimitiveKind { + // Lowest discriminant: at an equal order, a content-filter group-start is emitted before + // the group's own content so the renderer redirects rendering before any child draws. + FilterBoundaryStart, Shadow, #[default] Quad, @@ -195,6 +259,10 @@ pub(crate) enum PrimitiveKind { SubpixelSprite, PolychromeSprite, Surface, + BackdropFilter, + // Highest discriminant: at an equal order, a group-end is emitted after the group's content + // so the renderer composites the filtered group only once every child has been drawn. + FilterBoundaryEnd, } pub(crate) enum PaintOperation { @@ -214,6 +282,8 @@ pub enum Primitive { SubpixelSprite(SubpixelSprite), PolychromeSprite(PolychromeSprite), Surface(PaintSurface), + BackdropFilter(BackdropFilter), + FilterBoundary(FilterBoundary), } #[expect(missing_docs)] @@ -228,6 +298,8 @@ impl Primitive { Primitive::SubpixelSprite(sprite) => &sprite.bounds, Primitive::PolychromeSprite(sprite) => &sprite.bounds, Primitive::Surface(surface) => &surface.bounds, + Primitive::BackdropFilter(filter) => &filter.bounds, + Primitive::FilterBoundary(boundary) => &boundary.bounds, } } @@ -241,6 +313,8 @@ impl Primitive { Primitive::SubpixelSprite(sprite) => &sprite.content_mask, Primitive::PolychromeSprite(sprite) => &sprite.content_mask, Primitive::Surface(surface) => &surface.content_mask, + Primitive::BackdropFilter(filter) => &filter.content_mask, + Primitive::FilterBoundary(boundary) => &boundary.content_mask, } } } @@ -269,6 +343,10 @@ struct BatchIterator<'a> { polychrome_sprites_iter: Peekable>, surfaces_start: usize, surfaces_iter: Peekable>, + backdrop_filters_start: usize, + backdrop_filters_iter: Peekable>, + filter_boundaries_start: usize, + filter_boundaries_iter: Peekable>, } impl<'a> Iterator for BatchIterator<'a> { @@ -302,6 +380,20 @@ impl<'a> Iterator for BatchIterator<'a> { self.surfaces_iter.peek().map(|s| s.order), PrimitiveKind::Surface, ), + ( + self.backdrop_filters_iter.peek().map(|f| f.order), + PrimitiveKind::BackdropFilter, + ), + ( + self.filter_boundaries_iter.peek().map(|b| b.order), + // The same vec yields both start and end markers; the discriminant decides + // where the next marker sorts relative to draw batches at an equal order + // (start before content, end after). + match self.filter_boundaries_iter.peek() { + Some(boundary) if boundary.is_start => PrimitiveKind::FilterBoundaryStart, + _ => PrimitiveKind::FilterBoundaryEnd, + }, + ), ]; orders_and_kinds.sort_by_key(|(order, kind)| (order.unwrap_or(u32::MAX), *kind)); @@ -447,6 +539,30 @@ impl<'a> Iterator for BatchIterator<'a> { self.surfaces_start = surfaces_end; Some(PrimitiveBatch::Surfaces(surfaces_start..surfaces_end)) } + PrimitiveKind::BackdropFilter => { + let backdrop_filters_start = self.backdrop_filters_start; + let mut backdrop_filters_end = backdrop_filters_start + 1; + self.backdrop_filters_iter.next(); + while self + .backdrop_filters_iter + .next_if(|filter| (filter.order, batch_kind) < max_order_and_kind) + .is_some() + { + backdrop_filters_end += 1; + } + self.backdrop_filters_start = backdrop_filters_end; + Some(PrimitiveBatch::BackdropFilters( + backdrop_filters_start..backdrop_filters_end, + )) + } + // Boundaries are emitted one at a time (never merged) so the renderer can switch + // render targets at exactly the right point in the batch stream. + PrimitiveKind::FilterBoundaryStart | PrimitiveKind::FilterBoundaryEnd => { + let index = self.filter_boundaries_start; + self.filter_boundaries_iter.next(); + self.filter_boundaries_start = index + 1; + Some(PrimitiveBatch::FilterBoundary(index)) + } } } } @@ -479,6 +595,11 @@ pub enum PrimitiveBatch { range: Range, }, Surfaces(Range), + BackdropFilters(Range), + /// A single content-filter group boundary; index into [`Scene::filter_boundaries`]. Read + /// `is_start` to tell whether this opens the group (switch render target) or closes it + /// (filter the offscreen target and composite it back). + FilterBoundary(usize), } #[derive(Default, Debug, Copy, Clone)] @@ -543,6 +664,59 @@ impl From for Primitive { } } +/// A backdrop filter blurs (and may otherwise filter) the content already rendered behind +/// `bounds`, compositing the result into a rounded rectangle — the frosted-glass effect. +/// Emitted by [`crate::Window::paint_backdrop_filter`]; produces the CSS `backdrop-filter` effect. +#[derive(Default, Debug, Clone)] +#[expect(missing_docs)] +pub struct BackdropFilter { + pub order: DrawOrder, + pub bounds: Bounds, + pub content_mask: ContentMask, + pub corner_radii: Corners, + /// The filter chain applied to the backdrop, in scene (device-pixel) space. Identity filters + /// are dropped at paint time, so a `BackdropFilter` is only emitted when this is non-empty. + /// + /// Inline capacity is 4: a `SmallVec<[ScaledFilter; 4]>` is the same size as capacity 1 here + /// (the heap repr already occupies that space), so chains up to 4 filters avoid allocating + /// at no extra struct size. + pub filters: SmallVec<[ScaledFilter; 4]>, + /// Element opacity captured at paint time, multiplied into the composited result. + pub opacity: f32, +} + +impl From for Primitive { + fn from(filter: BackdropFilter) -> Self { + Primitive::BackdropFilter(filter) + } +} + +/// The start or end marker of a content-filter (`filter`) isolation group. The element's +/// subtree is painted between a matched start/end pair; the renderer redirects that span into +/// an offscreen target, filters it, and composites it back at `bounds`. Produces the CSS +/// `filter` effect (e.g. blurring the element and its children as a single group). +#[derive(Debug, Clone)] +#[expect(missing_docs)] +pub struct FilterBoundary { + pub order: DrawOrder, + pub bounds: Bounds, + pub content_mask: ContentMask, + pub corner_radii: Corners, + /// The filter chain applied to the isolated group, in scene (device-pixel) space. Identity + /// filters are dropped at paint time, so a `FilterBoundary` is only emitted when non-empty. + /// Inline capacity 4 (same struct size as 1 here — see [`BackdropFilter::filters`]). + pub filters: SmallVec<[ScaledFilter; 4]>, + pub opacity: f32, + /// `true` for the start marker (opens the group), `false` for the end marker (closes it). + pub is_start: bool, +} + +impl From for Primitive { + fn from(boundary: FilterBoundary) -> Self { + Primitive::FilterBoundary(boundary) + } +} + /// The style of a border. #[derive(Default, Debug, Clone, Copy, PartialEq, Eq, Hash, Serialize, Deserialize, JsonSchema)] #[repr(C)] @@ -903,3 +1077,171 @@ impl PathVertex { } } } + +#[cfg(test)] +mod tests { + use super::*; + use crate::{Point, Size}; + + fn sp(value: f32) -> ScaledPixels { + ScaledPixels(value) + } + + /// All test primitives cover the same region so the bounds tree assigns strictly + /// increasing orders in insertion order — making the expected batch order deterministic. + fn full_bounds() -> Bounds { + Bounds { + origin: Point { + x: sp(0.0), + y: sp(0.0), + }, + size: Size { + width: sp(100.0), + height: sp(100.0), + }, + } + } + + fn mask() -> ContentMask { + ContentMask { + bounds: full_bounds(), + } + } + + fn quad() -> Quad { + Quad { + bounds: full_bounds(), + content_mask: mask(), + ..Default::default() + } + } + + /// A 100x100 quad whose bounds don't overlap `full_bounds()` (used to exercise the + /// order-reuse path: non-overlapping content reuses low draw-orders). + fn detached_quad() -> Quad { + let bounds = Bounds { + origin: Point { + x: sp(200.0), + y: sp(200.0), + }, + size: Size { + width: sp(100.0), + height: sp(100.0), + }, + }; + Quad { + bounds, + content_mask: ContentMask { bounds }, + ..Default::default() + } + } + + fn boundary(is_start: bool) -> FilterBoundary { + FilterBoundary { + order: 0, + bounds: full_bounds(), + content_mask: mask(), + corner_radii: Corners::default(), + filters: smallvec::smallvec![ScaledFilter::Blur(sp(8.0))], + opacity: 1.0, + is_start, + } + } + + fn backdrop() -> BackdropFilter { + BackdropFilter { + bounds: full_bounds(), + content_mask: mask(), + corner_radii: Corners::default(), + filters: smallvec::smallvec![ScaledFilter::Blur(sp(20.0))], + opacity: 1.0, + ..Default::default() + } + } + + fn batch_kinds(scene: &mut Scene) -> Vec<&'static str> { + scene.finish(); + scene + .batches() + .map(|batch| match batch { + PrimitiveBatch::Quads(_) => "quad", + PrimitiveBatch::BackdropFilters(_) => "backdrop", + PrimitiveBatch::FilterBoundary(ix) => { + if scene.filter_boundaries[ix].is_start { + "start" + } else { + "end" + } + } + _ => "other", + }) + .collect() + } + + #[test] + fn content_filter_group_brackets_its_children() { + let mut scene = Scene::default(); + // Background painted before the filtered element. + scene.insert_primitive(quad()); + // A content-filtered element: start marker, its child, end marker. + scene.insert_primitive(boundary(true)); + scene.insert_primitive(quad()); + scene.insert_primitive(boundary(false)); + + // The start must precede the group's child and the end must follow it, so the + // renderer can redirect rendering for exactly the group's span. + assert_eq!( + batch_kinds(&mut scene), + vec!["quad", "start", "quad", "end"] + ); + } + + // Note: this validates only the *scene ordering* of nested filter boundaries (start/child/ + // end interleaving), not that a renderer actually isolates both levels — that depends on the + // backend's group-texture pool (see MAX_FILTER_DEPTH) and is exercised by the `blur` example. + #[test] + fn nested_content_filters_emit_well_nested_ordering() { + let mut scene = Scene::default(); + scene.insert_primitive(boundary(true)); // outer start + scene.insert_primitive(quad()); // outer child + scene.insert_primitive(boundary(true)); // inner start + scene.insert_primitive(quad()); // inner child + scene.insert_primitive(boundary(false)); // inner end + scene.insert_primitive(boundary(false)); // outer end + + assert_eq!( + batch_kinds(&mut scene), + vec!["start", "quad", "start", "quad", "end", "end"] + ); + } + + #[test] + fn content_after_a_filter_group_sorts_above_it() { + let mut scene = Scene::default(); + // A content-filtered element: start marker, its child, end marker. + scene.insert_primitive(boundary(true)); + scene.insert_primitive(quad()); + scene.insert_primitive(boundary(false)); + // A sibling painted after the group that does NOT overlap it. Without the close-time + // order-floor it would reuse the lowest order, tie with the start marker, and be swept + // into the group (start, quad, quad, end); it must instead sort after the end marker. + scene.insert_primitive(detached_quad()); + + assert_eq!( + batch_kinds(&mut scene), + vec!["start", "quad", "end", "quad"] + ); + } + + #[test] + fn backdrop_filter_sorts_before_a_later_overlapping_quad() { + let mut scene = Scene::default(); + // Content behind the frosted panel. + scene.insert_primitive(quad()); + // The panel: its backdrop snapshot, then its (translucent) background quad on top. + scene.insert_primitive(backdrop()); + scene.insert_primitive(quad()); + + assert_eq!(batch_kinds(&mut scene), vec!["quad", "backdrop", "quad"]); + } +} diff --git a/crates/gpui/src/style.rs b/crates/gpui/src/style.rs index 54e09bc37d..c66768d4b8 100644 --- a/crates/gpui/src/style.rs +++ b/crates/gpui/src/style.rs @@ -8,7 +8,8 @@ use crate::{ AbsoluteLength, App, Background, BackgroundTag, BorderStyle, Bounds, ContentMask, Corners, CornersRefinement, CursorStyle, DefiniteLength, DevicePixels, Edges, EdgesRefinement, Font, FontFallbacks, FontFeatures, FontStyle, FontWeight, GridLocation, Hsla, Length, Pixels, Point, - PointRefinement, Rgba, SharedString, Size, SizeRefinement, Styled, TextRun, Window, black, phi, + PointRefinement, Rgba, ScaledPixels, SharedString, Size, SizeRefinement, Styled, TextRun, + Window, black, phi, point, quad, rems, size, }; use collections::HashSet; @@ -287,6 +288,12 @@ pub struct Style { /// Box shadow of the element pub box_shadow: Vec, + /// Filters applied to this element's own content and children (CSS `filter`). + pub filter: Vec, + + /// Filters applied to the content rendered behind this element (CSS `backdrop-filter`). + pub backdrop_filter: Vec, + /// The text style of this element #[refineable] pub text: TextStyleRefinement, @@ -355,6 +362,50 @@ pub struct BoxShadow { pub inset: bool, } +/// A graphical filter that can be applied either to an element's own content +/// (via [`Styled::filter`], like CSS `filter`) or to the content rendered behind +/// it (via [`Styled::backdrop_filter`], like CSS `backdrop-filter`). +#[derive(Clone, Copy, Debug, PartialEq, Serialize, Deserialize, JsonSchema)] +pub enum Filter { + /// A gaussian blur with the given radius, in logical pixels. Maps to CSS `blur()`. + Blur(Pixels), +} + +impl Filter { + /// Whether this filter has no visible effect, so painting can skip it entirely (and the + /// element can avoid the offscreen isolation pass when *all* of its filters are identities). + /// + /// Each variant declares its own no-op case here rather than the pipeline special-casing + /// blur — adding a filter that this returns `true` for is silently dropped before it ever + /// reaches the renderer. + pub fn is_identity(&self) -> bool { + match self { + Filter::Blur(radius) => *radius <= Pixels::ZERO, + } + } + + /// Lower this logical-pixel filter into its scene-space ([`ScaledFilter`]) form for the + /// renderer, scaling any pixel magnitudes by `factor` (the window scale factor). + pub fn scale(&self, factor: f32) -> ScaledFilter { + match self { + Filter::Blur(radius) => ScaledFilter::Blur(radius.scale(factor)), + } + } +} + +/// The scene-space (device-pixel) form of a [`Filter`], carried on the scene primitives that the +/// renderers consume. Produced by [`Filter::scale`]; pixel magnitudes are in [`ScaledPixels`]. +/// +/// This is intentionally a separate enum from [`Filter`] (rather than reusing it) so the scene +/// stays in device space like every other primitive, and so the renderers `match` on it +/// exhaustively — adding a filter variant breaks each backend's match, forcing a deliberate +/// implement-or-decline decision per backend instead of silently rendering nothing. +#[derive(Clone, Copy, Debug, PartialEq)] +pub enum ScaledFilter { + /// A gaussian blur with the given radius, in scaled (device) pixels. + Blur(ScaledPixels), +} + /// How to handle whitespace in text #[derive(Copy, Clone, Debug, Default, PartialEq, Eq, Serialize, Deserialize, JsonSchema)] pub enum WhiteSpace { @@ -669,49 +720,68 @@ impl Style { window.paint_drop_shadows(bounds, corner_radii, &self.box_shadow); - let background_color = self.background.as_ref().and_then(Fill::color); - if background_color.is_some_and(|color| !color.is_transparent()) { - let mut border_color = match background_color { - Some(color) => match color.tag { - BackgroundTag::Solid - | BackgroundTag::PatternSlash - | BackgroundTag::Checkerboard => color.solid, - - BackgroundTag::LinearGradient => color - .colors - .first() - .map(|stop| stop.color) - .unwrap_or_default(), - }, - None => Hsla::default(), - }; - border_color.a = 0.; - window.paint_quad(quad( - bounds, - corner_radii, - background_color.unwrap_or_default(), - Edges::default(), - border_color, - self.border_style, - )); + // Blur the content behind this element before its (typically translucent) background + // is painted on top, so the background tints the frosted backdrop (CSS `backdrop-filter`). + if !self.backdrop_filter.is_empty() { + window.paint_backdrop_filter(bounds, corner_radii, &self.backdrop_filter); } - window.paint_inset_shadows(bounds, corner_radii, &self.box_shadow); + // The element's own box — background, inset shadows, children, and border — painted as a + // unit. A `filter` (CSS `filter`) wraps this whole unit so the renderer blurs the element + // and its children together as one group; without a filter it paints directly. + let paint_box = |window: &mut Window, cx: &mut App| { + let background_color = self.background.as_ref().and_then(Fill::color); + if background_color.is_some_and(|color| !color.is_transparent()) { + let mut border_color = match background_color { + Some(color) => match color.tag { + BackgroundTag::Solid + | BackgroundTag::PatternSlash + | BackgroundTag::Checkerboard => color.solid, - continuation(window, cx); + BackgroundTag::LinearGradient => color + .colors + .first() + .map(|stop| stop.color) + .unwrap_or_default(), + }, + None => Hsla::default(), + }; + border_color.a = 0.; + window.paint_quad(quad( + bounds, + corner_radii, + background_color.unwrap_or_default(), + Edges::default(), + border_color, + self.border_style, + )); + } - if self.is_border_visible() { - let border_widths = self.border_widths.to_pixels(rem_size); - let mut background = self.border_color.unwrap_or_default(); - background.a = 0.; - window.paint_quad(quad( - bounds, - corner_radii, - background, - border_widths, - self.border_color.unwrap_or_default(), - self.border_style, - )); + window.paint_inset_shadows(bounds, corner_radii, &self.box_shadow); + + continuation(window, cx); + + if self.is_border_visible() { + let border_widths = self.border_widths.to_pixels(rem_size); + let mut background = self.border_color.unwrap_or_default(); + background.a = 0.; + window.paint_quad(quad( + bounds, + corner_radii, + background, + border_widths, + self.border_color.unwrap_or_default(), + self.border_style, + )); + } + }; + + if self.filter.is_empty() { + paint_box(window, cx); + } else { + window.with_filter_layer(bounds, corner_radii, &self.filter, |window| { + paint_box(window, cx); + }); } #[cfg(debug_assertions)] @@ -765,6 +835,8 @@ impl Default for Style { border_style: BorderStyle::default(), corner_radii: Corners::default(), box_shadow: Default::default(), + filter: Default::default(), + backdrop_filter: Default::default(), text: TextStyleRefinement::default(), mouse_cursor: None, opacity: None, diff --git a/crates/gpui/src/styled.rs b/crates/gpui/src/styled.rs index e090ba973f..d95d444186 100644 --- a/crates/gpui/src/styled.rs +++ b/crates/gpui/src/styled.rs @@ -1,7 +1,7 @@ use crate::{ self as gpui, AbsoluteLength, AlignContent, AlignItems, AlignSelf, BorderStyle, CursorStyle, - DefiniteLength, Display, Fill, FlexDirection, FlexWrap, Font, FontFeatures, FontStyle, - FontWeight, GridPlacement, GridTemplate, Hsla, JustifyContent, Length, SharedString, + DefiniteLength, Display, Fill, Filter, FlexDirection, FlexWrap, Font, FontFeatures, FontStyle, + FontWeight, GridPlacement, GridTemplate, Hsla, JustifyContent, Length, Pixels, SharedString, StrikethroughStyle, StyleRefinement, TemplateColumnMinSize, TextAlign, TextOverflow, TextStyleRefinement, UnderlineStyle, WhiteSpace, px, relative, rems, }; @@ -33,6 +33,53 @@ pub trait Styled: Sized { gpui_macros::border_style_methods!(); gpui_macros::box_shadow_style_methods!(); + /// Blur this element's own content and children, like CSS `filter: blur()`. + /// + /// This isolates the element's subtree, blurs it as a group, and composites the + /// result back. To blur the content *behind* the element instead (frosted glass), + /// use [`Styled::backdrop_blur`]. + /// + /// *Appends* to the element's filter chain, so it composes with other convenience + /// setters (`.blur(8.).()`). To replace the whole chain, use + /// [`Styled::filter`]. + fn blur(mut self, radius: impl Into) -> Self { + self.style() + .filter + .get_or_insert_with(Vec::new) + .push(Filter::Blur(radius.into())); + self + } + + /// Set (replacing any existing) the full list of filters applied to this element's own + /// content, like CSS `filter`. To add a single filter to the chain instead, use the + /// convenience setters such as [`Styled::blur`]. + fn filter(mut self, filters: impl Into>) -> Self { + self.style().filter = Some(filters.into()); + self + } + + /// Blur the content rendered behind this element — a frosted-glass effect — like CSS + /// `backdrop-filter: blur()`. Typically paired with a translucent [`Styled::bg`] + /// so the background tints the blurred backdrop. + /// + /// *Appends* to the element's backdrop-filter chain. To replace the whole chain, use + /// [`Styled::backdrop_filter`]. + fn backdrop_blur(mut self, radius: impl Into) -> Self { + self.style() + .backdrop_filter + .get_or_insert_with(Vec::new) + .push(Filter::Blur(radius.into())); + self + } + + /// Set (replacing any existing) the full list of filters applied to the content behind this + /// element, like CSS `backdrop-filter`. To add a single filter to the chain instead, use the + /// convenience setters such as [`Styled::backdrop_blur`]. + fn backdrop_filter(mut self, filters: impl Into>) -> Self { + self.style().backdrop_filter = Some(filters.into()); + self + } + /// Sets the display type of the element to `block`. /// [Docs](https://tailwindcss.com/docs/display) fn block(mut self) -> Self { diff --git a/crates/gpui/src/window.rs b/crates/gpui/src/window.rs index b8805a3c4f..69fc04f41d 100644 --- a/crates/gpui/src/window.rs +++ b/crates/gpui/src/window.rs @@ -2,17 +2,18 @@ use crate::Inspector; use crate::{ Action, AnyDrag, AnyElement, AnyImageCache, AnyTooltip, AnyView, App, AppContext, Arena, Asset, - AsyncWindowContext, AvailableSpace, Background, BorderStyle, Bounds, BoxShadow, Capslock, - Context, Corners, CursorHideMode, CursorStyle, Decorations, DevicePixels, + AsyncWindowContext, AvailableSpace, BackdropFilter, Background, BorderStyle, Bounds, BoxShadow, + Capslock, Context, Corners, CursorHideMode, CursorStyle, Decorations, DevicePixels, DispatchActionListener, DispatchNodeId, DispatchTree, DisplayId, Edges, Effect, Entity, - EntityId, EventEmitter, FileDropEvent, FontId, Global, GlobalElementId, GlyphId, GpuSpecs, - Hsla, InputHandler, IsZero, KeyBinding, KeyContext, KeyDownEvent, KeyEvent, Keystroke, - KeystrokeEvent, LayoutId, Lerp, LineLayoutIndex, Modifiers, ModifiersChangedEvent, + EntityId, EventEmitter, FileDropEvent, Filter, FilterBoundary, FontId, Global, GlobalElementId, + GlyphId, GpuSpecs, Hsla, InputHandler, IsZero, KeyBinding, KeyContext, KeyDownEvent, KeyEvent, + Keystroke, KeystrokeEvent, LayoutId, Lerp, LineLayoutIndex, Modifiers, ModifiersChangedEvent, MonochromeSprite, MouseButton, MouseEvent, MouseMoveEvent, MouseUpEvent, Path, Pixels, PlatformAtlas, PlatformDisplay, PlatformInput, PlatformInputHandler, PlatformWindow, Point, PolychromeSprite, Priority, PromptButton, PromptLevel, Quad, Render, RenderGlyphParams, RenderImage, RenderImageParams, RenderSvgParams, Replay, ResizeEdge, SMOOTH_SVG_SCALE_FACTOR, - SUBPIXEL_VARIANTS_X, SUBPIXEL_VARIANTS_Y, ScaledPixels, Scene, Shadow, SharedString, Size, + SUBPIXEL_VARIANTS_X, SUBPIXEL_VARIANTS_Y, ScaledFilter, ScaledPixels, Scene, Shadow, + SharedString, Size, StrikethroughStyle, Style, SubpixelSprite, SubscriberSet, Subscription, SystemWindowTab, SystemWindowTabController, TabStopMap, TaffyLayoutEngine, Task, TextRenderingMode, TextStyle, TextStyleRefinement, ThermalState, TransformationMatrix, Transition, TransitionState, @@ -2938,6 +2939,11 @@ impl Window { return; } + // Deferred draws are overlays (tooltips, popovers, drag images) and must sort above the + // whole main scene. Raise the order floor so they do — this also keeps a deferred + // backdrop's order from falling inside a content-filter order range left by the main scene. + self.next_frame.scene.raise_order_floor(); + let traversal_order = self.deferred_draw_traversal_order(); let mut deferred_draws = mem::take(&mut self.next_frame.deferred_draws); for deferred_draw_ix in traversal_order { @@ -3677,6 +3683,98 @@ impl Window { } } + /// Paint a backdrop filter into the scene for the next frame at the current z-index. The + /// renderer blurs the content already painted behind `bounds` and composites the result + /// into the rounded rectangle described by `bounds` and `corner_radii` — the CSS + /// `backdrop-filter` effect (frosted glass). Typically the element then paints a translucent + /// background quad on top so its color tints the blurred backdrop. + /// + /// Does nothing when `filters` produce no visible blur. + /// + /// This method should only be called as part of the paint phase of element drawing. + pub fn paint_backdrop_filter( + &mut self, + bounds: Bounds, + corner_radii: Corners, + filters: &[Filter], + ) { + self.invalidator.debug_assert_paint(); + + let scale_factor = self.scale_factor(); + let filters: SmallVec<[ScaledFilter; 4]> = filters + .iter() + .filter(|filter| !filter.is_identity()) + .map(|filter| filter.scale(scale_factor)) + .collect(); + if filters.is_empty() { + return; + } + + self.next_frame.scene.insert_primitive(BackdropFilter { + order: 0, + bounds: self.snap_bounds(bounds), + content_mask: self.snapped_content_mask(), + corner_radii: corner_radii.scale(scale_factor), + filters, + opacity: self.element_opacity(), + }); + } + + /// Isolate the painting performed by `f` into a content-filter group: the renderer renders + /// everything `f` paints into an offscreen target, blurs it as a single layer, and + /// composites the result back into the rounded rectangle described by `bounds` and + /// `corner_radii` — the CSS `filter` effect (e.g. blurring an element and its children). + /// + /// When `filters` produce no visible blur this simply runs `f` with no offscreen + /// indirection. + /// + /// This method should only be called as part of the paint phase of element drawing. + pub fn with_filter_layer( + &mut self, + bounds: Bounds, + corner_radii: Corners, + filters: &[Filter], + f: impl FnOnce(&mut Self) -> R, + ) -> R { + self.invalidator.debug_assert_paint(); + + let scale_factor = self.scale_factor(); + let filters: SmallVec<[ScaledFilter; 4]> = filters + .iter() + .filter(|filter| !filter.is_identity()) + .map(|filter| filter.scale(scale_factor)) + .collect(); + if filters.is_empty() { + return f(self); + } + + // Snapshot the (scaled) group parameters once so the start and end markers agree. + // + // `opacity` is 1.0 — NOT `element_opacity()`. The group's children/bg/border are painted + // through the normal paint methods while `element_opacity` is still in effect, so they + // already carry the element's opacity (consistent with gpui's per-primitive opacity for + // non-filtered elements). Re-applying it at composite time would double it (e.g. + // `.blur(r).opacity(0.5)` would render at 0.25 instead of 0.5). + let boundary = FilterBoundary { + order: 0, + bounds: self.snap_bounds(bounds), + content_mask: self.snapped_content_mask(), + corner_radii: corner_radii.scale(scale_factor), + filters, + opacity: 1.0, + is_start: true, + }; + + self.next_frame.scene.insert_primitive(boundary.clone()); + let result = f(self); + self.next_frame.scene.insert_primitive(FilterBoundary { + is_start: false, + ..boundary + }); + + result + } + /// Paint one or more quads into the scene for the next frame at the current stacking context. /// Quads are colored rectangular regions with an optional background, border, and corner radius. /// see [`fill`], [`outline`], and [`quad`] to construct this type. diff --git a/crates/gpui_macos/src/metal_renderer.rs b/crates/gpui_macos/src/metal_renderer.rs index 73b53ce6ea..0e2a408564 100644 --- a/crates/gpui_macos/src/metal_renderer.rs +++ b/crates/gpui_macos/src/metal_renderer.rs @@ -7,10 +7,22 @@ use cocoa::{ quartzcore::AutoresizingMask, }; use gpui::{ - AtlasTextureId, Background, Bounds, ContentMask, DevicePixels, MonochromeSprite, PaintSurface, - Path, Point, PolychromeSprite, PrimitiveBatch, Quad, ScaledPixels, Scene, Shadow, Size, - Surface, Underline, point, size, + AtlasTextureId, Background, Bounds, ContentMask, Corners, DevicePixels, FilterBoundary, + MonochromeSprite, PaintSurface, Path, Point, PolychromeSprite, PrimitiveBatch, Quad, + ScaledFilter, ScaledPixels, Scene, Shadow, Size, Surface, Underline, point, size, }; + +/// The largest blur radius in a scene-space filter chain, in device pixels — used to size the +/// blur kernel and the dilated region the blur passes are scissored to. +/// +/// The `match` is exhaustive on purpose: adding a [`ScaledFilter`] variant breaks it here, +/// forcing this backend to handle (or deliberately ignore) the new filter rather than silently +/// dropping it. +fn max_blur_radius(filters: &[ScaledFilter]) -> f32 { + filters.iter().fold(0.0, |acc, filter| match filter { + ScaledFilter::Blur(radius) => acc.max(radius.0), + }) +} #[cfg(any(test, feature = "test-support"))] use image::RgbaImage; @@ -40,6 +52,12 @@ const SHADERS_SOURCE_FILE: &str = include_str!(concat!(env!("OUT_DIR"), "/stitch // https://developer.apple.com/documentation/metal/mtldevice/1433355-supportstexturesamplecount const PATH_SAMPLE_COUNT: u32 = 4; +/// Number of content-filter (`filter`) nesting levels that get their own isolated group texture. +/// Two covers the realistic "a blurred element inside another blurred element" case; deeper nests +/// render inline (unblurred at the inner level) rather than allocating unbounded VRAM. Must match +/// the wgpu backend's `MAX_FILTER_DEPTH` so nested blur renders consistently across platforms. +const MAX_FILTER_DEPTH: usize = 2; + pub(crate) type Context = Arc>; pub(crate) type Renderer = MetalRenderer; @@ -125,6 +143,11 @@ pub(crate) struct MetalRenderer { monochrome_sprites_pipeline_state: metal::RenderPipelineState, polychrome_sprites_pipeline_state: metal::RenderPipelineState, surfaces_pipeline_state: metal::RenderPipelineState, + // Blur pipelines: downsample (no blend, also used for the final blit), separable gaussian + // (no blend), and composite (alpha blend into a rounded rect). See `shaders.metal`. + blur_downsample_pipeline_state: metal::RenderPipelineState, + blur_pipeline_state: metal::RenderPipelineState, + blur_composite_pipeline_state: metal::RenderPipelineState, unit_vertices: metal::Buffer, #[allow(clippy::arc_with_non_send_sync)] instance_buffer_pool: Arc>, @@ -132,9 +155,68 @@ pub(crate) struct MetalRenderer { core_video_texture_cache: core_video::metal_texture_cache::CVMetalTextureCache, path_intermediate_texture: Option, path_intermediate_msaa_texture: Option, + // Offscreen scene target (the scene is rendered here, then blitted to the drawable, so blur + // passes can sample already-painted content), the half-res ping/pong blur targets, and a + // full-res target for content-filter groups. + scene_color_texture: Option, + blur_ping_texture: Option, + blur_pong_texture: Option, + /// Full-resolution offscreen targets a content-filter (`filter`) group renders into before + /// being blurred and composited back. One per nesting level (indexed by isolation depth) so + /// nested content blurs isolate correctly, up to [`MAX_FILTER_DEPTH`]; deeper nests render + /// inline. + group_textures: Vec, path_sample_count: u32, } +/// Mirrors the `BlurParams` struct in `shaders.metal`. Passed to the blur pipelines via +/// `setVertexBytes`/`setFragmentBytes`. +#[repr(C)] +#[derive(Clone, Copy)] +struct BlurUniform { + bounds: Bounds, + content_mask: Bounds, + corner_radii: Corners, + direction: [f32; 2], + sigma: f32, + opacity: f32, + tap_count: f32, + /// 1.0 clips the composite to the rounded rect (backdrop); 0.0 lets content blur bleed past + /// its bounds like CSS `filter: blur`. + clip_rounded: f32, + /// 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the + /// origin, so a stationary element blurs identically at every window size); 0.0 = 1:1 copy + /// (the scene blit, which must not downsample). Downsample pass only. + downsample: f32, + /// Spacing between taps in pixels (gaussian passes only); >1 lets `tap_count` taps span very + /// large radii without truncating the gaussian, matching the wgpu backend. + tap_step: f32, +} + +impl Default for BlurUniform { + fn default() -> Self { + BlurUniform { + bounds: Bounds::default(), + content_mask: Bounds::default(), + corner_radii: Corners::default(), + direction: [0.0, 0.0], + sigma: 0.0, + opacity: 1.0, + tap_count: 0.0, + clip_rounded: 0.0, + downsample: 0.0, + tap_step: 0.0, + } + } +} + +#[repr(C)] +enum BlurInputIndex { + Vertices = 0, + Params = 1, + ViewportSize = 2, +} + #[repr(C)] pub struct PathRasterizationVertex { pub xy_position: Point, @@ -318,6 +400,32 @@ impl MetalRenderer { "surface_fragment", MTLPixelFormat::BGRA8Unorm, ); + let blur_downsample_pipeline_state = build_blur_pipeline_state( + &device, + &library, + "blur_downsample", + "blur_fullscreen_vertex", + "blur_downsample_fragment", + MTLPixelFormat::BGRA8Unorm, + ); + let blur_pipeline_state = build_blur_pipeline_state( + &device, + &library, + "blur", + "blur_fullscreen_vertex", + "blur_fragment", + MTLPixelFormat::BGRA8Unorm, + ); + // Premultiplied blend (One / OneMinusSourceAlpha) — the composite outputs a premultiplied + // blurred sample; straight-alpha blending would darken the faded edges. + let blur_composite_pipeline_state = build_path_sprite_pipeline_state( + &device, + &library, + "blur_composite", + "blur_composite_vertex", + "blur_composite_fragment", + MTLPixelFormat::BGRA8Unorm, + ); let command_queue = device.new_command_queue(); let sprite_atlas = Arc::new(MetalAtlas::new(device.clone(), is_apple_gpu)); @@ -340,12 +448,19 @@ impl MetalRenderer { monochrome_sprites_pipeline_state, polychrome_sprites_pipeline_state, surfaces_pipeline_state, + blur_downsample_pipeline_state, + blur_pipeline_state, + blur_composite_pipeline_state, unit_vertices, instance_buffer_pool, sprite_atlas, core_video_texture_cache, path_intermediate_texture: None, path_intermediate_msaa_texture: None, + scene_color_texture: None, + blur_ping_texture: None, + blur_pong_texture: None, + group_textures: Vec::new(), path_sample_count: PATH_SAMPLE_COUNT, } } @@ -395,6 +510,10 @@ impl MetalRenderer { if size.width.0 <= 0 || size.height.0 <= 0 { self.path_intermediate_texture = None; self.path_intermediate_msaa_texture = None; + self.scene_color_texture = None; + self.blur_ping_texture = None; + self.blur_pong_texture = None; + self.group_textures.clear(); return; } @@ -407,6 +526,27 @@ impl MetalRenderer { .set_usage(metal::MTLTextureUsage::RenderTarget | metal::MTLTextureUsage::ShaderRead); self.path_intermediate_texture = Some(self.device.new_texture(&texture_descriptor)); + // Full-res scene + group targets, and half-res ping/pong blur targets. + let make_color_texture = |width: u64, height: u64| { + let descriptor = metal::TextureDescriptor::new(); + descriptor.set_width(width.max(1)); + descriptor.set_height(height.max(1)); + descriptor.set_pixel_format(metal::MTLPixelFormat::BGRA8Unorm); + descriptor.set_storage_mode(metal::MTLStorageMode::Private); + descriptor.set_usage( + metal::MTLTextureUsage::RenderTarget | metal::MTLTextureUsage::ShaderRead, + ); + self.device.new_texture(&descriptor) + }; + let full_w = size.width.0 as u64; + let full_h = size.height.0 as u64; + self.scene_color_texture = Some(make_color_texture(full_w, full_h)); + self.group_textures = (0..MAX_FILTER_DEPTH) + .map(|_| make_color_texture(full_w, full_h)) + .collect(); + self.blur_ping_texture = Some(make_color_texture(full_w / 2, full_h / 2)); + self.blur_pong_texture = Some(make_color_texture(full_w / 2, full_h / 2)); + if self.path_sample_count > 1 { // https://developer.apple.com/documentation/metal/choosing-a-resource-storage-mode-for-apple-gpus // Rendering MSAA textures are done in a single pass, so we can use memory-less storage on Apple Silicon @@ -751,9 +891,31 @@ impl MetalRenderer { let alpha = if self.opaque { 1. } else { 0. }; let mut instance_offset = 0; + // Render the scene into an offscreen color texture (so filters can sample it), then + // blit it to `texture`. Owned clones keep the textures borrowable without borrowing + // `self` across the batch loop (which calls `&mut self` methods like `draw_surfaces`). + // Only route through the offscreen scene texture when the scene actually contains blur + // filters; otherwise render straight to `texture` exactly as before (no regression, no + // extra blit for the common case). + let use_offscreen = + !scene.backdrop_filters.is_empty() || !scene.filter_boundaries.is_empty(); + let scene_color_owned = self.scene_color_texture.clone(); + let blur_ping_owned = self.blur_ping_texture.clone(); + let blur_pong_owned = self.blur_pong_texture.clone(); + let group_owned = self.group_textures.clone(); + let scene_color: &metal::TextureRef = if use_offscreen { + scene_color_owned.as_deref().unwrap_or(texture) + } else { + texture + }; + // The active render target; switches to the group texture inside a content-filter group. + let mut current_target: &metal::TextureRef = scene_color; + // (boundary, parent target to composite back into, whether this level is isolated). + let mut filter_stack: Vec<(FilterBoundary, &metal::TextureRef, bool)> = Vec::new(); + let mut command_encoder = new_command_encoder_for_texture( command_buffer, - texture, + current_target, viewport_size, |color_attachment| { color_attachment.set_load_action(metal::MTLLoadAction::Clear); @@ -791,7 +953,7 @@ impl MetalRenderer { command_encoder = new_command_encoder_for_texture( command_buffer, - texture, + current_target, viewport_size, |color_attachment| { color_attachment.set_load_action(metal::MTLLoadAction::Load); @@ -842,6 +1004,98 @@ impl MetalRenderer { viewport_size, command_encoder, ), + PrimitiveBatch::BackdropFilters(range) => { + command_encoder.end_encoding(); + if let (Some(ping), Some(pong)) = + (blur_ping_owned.as_deref(), blur_pong_owned.as_deref()) + { + for filter in &scene.backdrop_filters[range] { + self.metal_blur_and_composite( + command_buffer, + current_target, + current_target, + ping, + pong, + viewport_size, + filter.bounds, + filter.content_mask.bounds, + filter.corner_radii, + max_blur_radius(&filter.filters), + filter.opacity, + true, + ); + } + } + command_encoder = new_command_encoder_for_texture( + command_buffer, + current_target, + viewport_size, + |color_attachment| { + color_attachment.set_load_action(metal::MTLLoadAction::Load); + }, + ); + true + } + PrimitiveBatch::FilterBoundary(ix) => { + let boundary = scene.filter_boundaries[ix].clone(); + if boundary.is_start { + // Each isolated nesting level uses its own group texture from the pool + // (indexed by current isolation depth). Beyond the pool size + // (MAX_FILTER_DEPTH) deeper filters render inline without isolation rather + // than corrupting an outer group. + let depth = filter_stack.iter().filter(|entry| entry.2).count(); + if depth < group_owned.len() { + command_encoder.end_encoding(); + let parent = current_target; + current_target = group_owned[depth].as_ref(); + filter_stack.push((boundary, parent, true)); + command_encoder = new_command_encoder_for_texture( + command_buffer, + current_target, + viewport_size, + |color_attachment| { + color_attachment.set_load_action(metal::MTLLoadAction::Clear); + color_attachment + .set_clear_color(metal::MTLClearColor::new(0., 0., 0., 0.)); + }, + ); + } else { + filter_stack.push((boundary, current_target, false)); + } + } else if let Some((boundary, parent, isolated)) = filter_stack.pop() { + if isolated { + command_encoder.end_encoding(); + if let (Some(ping), Some(pong)) = + (blur_ping_owned.as_deref(), blur_pong_owned.as_deref()) + { + self.metal_blur_and_composite( + command_buffer, + current_target, + parent, + ping, + pong, + viewport_size, + boundary.bounds, + boundary.content_mask.bounds, + boundary.corner_radii, + max_blur_radius(&boundary.filters), + boundary.opacity, + false, + ); + } + current_target = parent; + command_encoder = new_command_encoder_for_texture( + command_buffer, + current_target, + viewport_size, + |color_attachment| { + color_attachment.set_load_action(metal::MTLLoadAction::Load); + }, + ); + } + } + true + } PrimitiveBatch::SubpixelSprites { .. } => unreachable!(), }; if !ok { @@ -861,6 +1115,19 @@ impl MetalRenderer { command_encoder.end_encoding(); + // Present the offscreen scene by copying it into the drawable/target texture. + if use_offscreen && scene_color_owned.is_some() { + self.run_metal_blur_pass( + command_buffer, + &self.blur_downsample_pipeline_state, + texture, + scene_color, + viewport_size, + BlurUniform::default(), + false, + ); + } + if !self.is_unified_memory { // Sync the instance buffer to the GPU instance_buffer.metal_buffer.did_modify_range(NSRange { @@ -872,6 +1139,170 @@ impl MetalRenderer { Ok(command_buffer.to_owned()) } + /// Run a single blur pass: draw a full-screen (or composite) quad sampling `source` into + /// `target`. `params` is supplied to both shader stages; `load` keeps existing target + /// contents (used by the composite), otherwise the target is cleared. + #[allow(clippy::too_many_arguments)] + fn run_metal_blur_pass( + &self, + command_buffer: &metal::CommandBufferRef, + pipeline: &metal::RenderPipelineState, + target: &metal::TextureRef, + source: &metal::TextureRef, + target_viewport: Size, + params: BlurUniform, + load: bool, + ) { + let encoder = new_command_encoder_for_texture( + command_buffer, + target, + target_viewport, + |color_attachment| { + if load { + color_attachment.set_load_action(metal::MTLLoadAction::Load); + } else { + color_attachment.set_load_action(metal::MTLLoadAction::Clear); + color_attachment.set_clear_color(metal::MTLClearColor::new(0., 0., 0., 0.)); + } + }, + ); + encoder.set_render_pipeline_state(pipeline); + encoder.set_vertex_buffer( + BlurInputIndex::Vertices as u64, + Some(&self.unit_vertices), + 0, + ); + encoder.set_vertex_bytes( + BlurInputIndex::Params as u64, + mem::size_of::() as u64, + ¶ms as *const BlurUniform as *const _, + ); + encoder.set_vertex_bytes( + BlurInputIndex::ViewportSize as u64, + mem::size_of_val(&target_viewport) as u64, + &target_viewport as *const Size as *const _, + ); + encoder.set_fragment_bytes( + BlurInputIndex::Params as u64, + mem::size_of::() as u64, + ¶ms as *const BlurUniform as *const _, + ); + encoder.set_fragment_bytes( + BlurInputIndex::ViewportSize as u64, + mem::size_of_val(&target_viewport) as u64, + &target_viewport as *const Size as *const _, + ); + encoder.set_fragment_texture(0, Some(source)); + encoder.draw_primitives(metal::MTLPrimitiveType::Triangle, 0, 6); + encoder.end_encoding(); + } + + /// Blur `source` (full-resolution) using the half-res ping/pong textures and composite the + /// result into `target`, clipped to `bounds`/`corner_radii`/`content_mask` and modulated by + /// `opacity`. Shared by the backdrop and content-filter paths. + #[allow(clippy::too_many_arguments)] + fn metal_blur_and_composite( + &self, + command_buffer: &metal::CommandBufferRef, + source: &metal::TextureRef, + target: &metal::TextureRef, + ping: &metal::TextureRef, + pong: &metal::TextureRef, + viewport_size: Size, + bounds: Bounds, + content_mask: Bounds, + corner_radii: Corners, + blur_radius: f32, + opacity: f32, + // Backdrop clips to the rounded rect; content (`filter`) bleeds past its bounds. + clip_rounded: bool, + ) { + // Sigma is halved because the blur runs at half resolution. + let sigma = (blur_radius * 0.5).max(0.0); + if sigma <= 0.0 { + return; + } + // Span ±3σ. If that needs more than 32 taps, spread the taps apart (tap_step > 1) rather + // than truncating the kernel — keeps very large radii from clipping. Matches wgpu. + let ideal_taps = (3.0 * sigma).ceil(); + let tap_count = ideal_taps.clamp(1.0, 32.0); + let tap_step = (ideal_taps / tap_count).max(1.0); + // Content blur bleeds ~3·radius past the box, so its composite quad covers a dilated rect. + let composite_bounds = if clip_rounded { + bounds + } else { + bounds.dilate(ScaledPixels(3.0 * blur_radius)) + }; + let half = Size { + width: DevicePixels((i32::from(viewport_size.width) / 2).max(1)), + height: DevicePixels((i32::from(viewport_size.height) / 2).max(1)), + }; + let half_w = i32::from(half.width) as f32; + let half_h = i32::from(half.height) as f32; + + // Downsample source -> ping, then separable gaussian ping -> pong -> ping. + self.run_metal_blur_pass( + command_buffer, + &self.blur_downsample_pipeline_state, + ping, + source, + half, + BlurUniform { + downsample: 1.0, + ..Default::default() + }, + false, + ); + self.run_metal_blur_pass( + command_buffer, + &self.blur_pipeline_state, + pong, + ping, + half, + BlurUniform { + direction: [1.0 / half_w, 0.0], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + false, + ); + self.run_metal_blur_pass( + command_buffer, + &self.blur_pipeline_state, + ping, + pong, + half, + BlurUniform { + direction: [0.0, 1.0 / half_h], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + false, + ); + + // Composite the blurred result into the target (preserving its contents). + self.run_metal_blur_pass( + command_buffer, + &self.blur_composite_pipeline_state, + target, + ping, + viewport_size, + BlurUniform { + bounds: composite_bounds, + content_mask, + corner_radii, + opacity, + clip_rounded: if clip_rounded { 1.0 } else { 0.0 }, + ..Default::default() + }, + true, + ); + } + fn draw_paths_to_intermediate( &self, paths: &[Path], @@ -1615,6 +2046,36 @@ fn build_path_rasterization_pipeline_state( .expect("could not create render pipeline state") } +// Blur downsample/gaussian passes overwrite their target (no blending). The composite pass +// uses the normal alpha-blending pipeline (`build_pipeline_state`) instead. +fn build_blur_pipeline_state( + device: &metal::DeviceRef, + library: &metal::LibraryRef, + label: &str, + vertex_fn_name: &str, + fragment_fn_name: &str, + pixel_format: metal::MTLPixelFormat, +) -> metal::RenderPipelineState { + let vertex_fn = library + .get_function(vertex_fn_name, None) + .expect("error locating vertex function"); + let fragment_fn = library + .get_function(fragment_fn_name, None) + .expect("error locating fragment function"); + + let descriptor = metal::RenderPipelineDescriptor::new(); + descriptor.set_label(label); + descriptor.set_vertex_function(Some(vertex_fn.as_ref())); + descriptor.set_fragment_function(Some(fragment_fn.as_ref())); + let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + color_attachment.set_pixel_format(pixel_format); + color_attachment.set_blending_enabled(false); + + device + .new_render_pipeline_state(&descriptor) + .expect("could not create render pipeline state") +} + // Align to multiples of 256 make Metal happy. fn align_offset(offset: &mut usize) { *offset = (*offset).div_ceil(256) * 256; diff --git a/crates/gpui_macos/src/shaders.metal b/crates/gpui_macos/src/shaders.metal index 2b52bb9ecc..155ceae162 100644 --- a/crates/gpui_macos/src/shaders.metal +++ b/crates/gpui_macos/src/shaders.metal @@ -1277,3 +1277,133 @@ float4 fill_color(Background background, return color; } + +// --- blur --- // +// +// Shared by backdrop (`backdrop-filter`) and content (`filter`) blur. Three passes: +// downsample (full -> half res), separable gaussian (run twice), and a composite that +// samples the blurred texture into a rounded rectangle. `BlurParams` is supplied via +// `setFragmentBytes`/`setVertexBytes` and mirrors the Rust `BlurUniform` struct exactly. +// +// Buffer/texture indices (raw, matching gpui_macos::metal_renderer::BlurInputIndex): +// buffer(0) = unit vertices, buffer(1) = BlurParams, buffer(2) = viewport size +// texture(0) = source + +struct BlurParams { + Bounds_ScaledPixels bounds; + Bounds_ScaledPixels content_mask; + Corners_ScaledPixels corner_radii; + float2 direction; + float sigma; + float opacity; + float tap_count; + float clip_rounded; + // 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the origin + // so a stationary element blurs identically at every window size); 0.0 = 1:1 copy (scene blit). + float downsample; + // Spacing between taps in pixels (gaussian passes only); >1 lets `tap_count` taps span very + // large radii without truncating the gaussian. + float tap_step; +}; + +struct BlurVertexOutput { + float4 position [[position]]; + float2 uv; +}; + +vertex BlurVertexOutput blur_fullscreen_vertex( + uint unit_vertex_id [[vertex_id]], + constant float2 *unit_vertices [[buffer(0)]]) { + float2 uv = unit_vertices[unit_vertex_id]; + BlurVertexOutput out; + out.position = float4(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, 0.0, 1.0); + out.uv = uv; + return out; +} + +fragment float4 blur_downsample_fragment( + BlurVertexOutput input [[stage_in]], + texture2d source [[texture(0)]], + constant BlurParams ¶ms [[buffer(1)]]) { + constexpr sampler s(mag_filter::linear, min_filter::linear); + if (params.downsample > 0.5) { + // Snapped 2:1 box downsample. Half-res texel `px` samples source at full-res coordinate + // 2*px + 1 (the boundary between source texels 2*px and 2*px+1), so one bilinear tap averages + // exactly that pair. Anchored to the origin and independent of the viewport size, so an element + // at fixed pixels blurs identically at every window size — otherwise the implicit floor(W/2) + // grid stretches and the halo wobbles by ~1px on resize. Uses the source's own dimensions + // because this pass binds the half-res target size as the viewport. + float2 dst = floor(input.position.xy); + float2 src_size = float2(float(source.get_width()), float(source.get_height())); + float2 src_uv = (dst * 2.0 + 1.0) / src_size; + return source.sample(s, src_uv); + } + // 1:1 copy at matching resolution (used to blit the offscreen scene into the drawable). + return source.sample(s, input.uv); +} + +fragment float4 blur_fragment( + BlurVertexOutput input [[stage_in]], + texture2d source [[texture(0)]], + constant BlurParams ¶ms [[buffer(1)]]) { + constexpr sampler s(mag_filter::linear, min_filter::linear); + int taps = int(params.tap_count); + float4 color = float4(0.0); + float weight_sum = 0.0; + for (int i = -taps; i <= taps; i++) { + float offset = float(i) * params.tap_step; + float w = gaussian(offset, params.sigma); + color += source.sample(s, input.uv + params.direction * offset) * w; + weight_sum += w; + } + return color / max(weight_sum, 1e-5); +} + +struct BlurCompositeVertexOutput { + float4 position [[position]]; + float clip_distance [[clip_distance]][4]; +}; + +struct BlurCompositeFragmentInput { + float4 position [[position]]; +}; + +vertex BlurCompositeVertexOutput blur_composite_vertex( + uint unit_vertex_id [[vertex_id]], + constant float2 *unit_vertices [[buffer(0)]], + constant BlurParams ¶ms [[buffer(1)]], + constant Size_DevicePixels *viewport_size [[buffer(2)]]) { + float2 unit_vertex = unit_vertices[unit_vertex_id]; + BlurCompositeVertexOutput out; + out.position = to_device_position(unit_vertex, params.bounds, viewport_size); + float4 clip = distance_from_clip_rect(unit_vertex, params.bounds, params.content_mask); + out.clip_distance[0] = clip.x; + out.clip_distance[1] = clip.y; + out.clip_distance[2] = clip.z; + out.clip_distance[3] = clip.w; + return out; +} + +fragment float4 blur_composite_fragment( + BlurCompositeFragmentInput input [[stage_in]], + texture2d source [[texture(0)]], + constant BlurParams ¶ms [[buffer(1)]], + constant Size_DevicePixels *viewport_size [[buffer(2)]]) { + constexpr sampler s(mag_filter::linear, min_filter::linear); + // Sample the half-res blur by screen position, on the SAME fixed 2:1 grid the snapped downsample + // wrote (anchored at the origin, independent of viewport parity): 2 * the half-res texture size + // maps screen pixel p to half-res texel p/2 at every window size, so it doesn't wobble on resize. + float2 half_size = float2(float(source.get_width()), float(source.get_height())); + float2 uv = input.position.xy / (2.0 * half_size); + float4 blurred = source.sample(s, uv); + // Backdrop clips to the rounded rect (the panel has a defined shape); content blur bleeds past + // its bounds like CSS `filter: blur`, so its shape comes from the blurred group's own alpha. + float dist = quad_sdf(input.position.xy, params.bounds, params.corner_radii); + float coverage = params.clip_rounded > 0.5 ? saturate(0.5 - dist) : 1.0; + // The blurred sample is premultiplied (blurring against the transparent surround scales rgb with + // the fading alpha), so output premultiplied and use a premultiplied-blend pipeline. A backdrop's + // scene is opaque (so this replaces); a content-filter group is transparent outside its subtree + // (so the target shows through there instead of darkening). + float a = coverage * params.opacity; + return float4(blurred.rgb * a, blurred.a * a); +} diff --git a/crates/gpui_wgpu/src/shaders.wgsl b/crates/gpui_wgpu/src/shaders.wgsl index b318a6829c..e9a5c46a98 100644 --- a/crates/gpui_wgpu/src/shaders.wgsl +++ b/crates/gpui_wgpu/src/shaders.wgsl @@ -1348,3 +1348,136 @@ fn fs_surface(input: SurfaceVarying) -> @location(0) vec4 { return textureSampleLevel(t_surface, s_surface, input.texture_position, 0.0); } + +// --- blur --- // +// +// Backdrop and content filters share these passes: +// 1. `fs_blur_downsample` copies a source texture into the half-resolution blur texture +// (also reused to blit the offscreen scene into the swapchain). +// 2. `fs_blur` runs one axis of a separable gaussian; the host invokes it twice. +// 3. `fs_blur_composite` samples the blurred texture and composites it into a rounded +// rectangle, clipped and modulated by opacity. +// +// Notes (review #5, #7): the scene/blur textures use the swapchain's (typically non-sRGB) +// format, so the gaussian runs on gamma-encoded values rather than linear light — consistent +// with the rest of gpui's compositing and close to what browsers do; bright detail darkens +// slightly. A content-filter group's texture is transparent outside the painted subtree, so +// the blur bleeds toward transparent at the group's edges (a soft edge ring) before the +// rounded-rect clip — this matches CSS `filter: blur` edge behaviour. + +struct BlurParams { + bounds: Bounds, + content_mask: Bounds, + corner_radii: vec4, + direction: vec2, + sigma: f32, + opacity: f32, + tap_count: f32, + // Spacing between taps, in pixels. >1 when the radius is so large the kernel would need more + // than `tap_count` taps to span ±3σ — the taps spread out instead of truncating the gaussian. + tap_step: f32, + // 1.0 = clip the composite to the rounded rect (backdrop); 0.0 = let the blurred result fade + // out on its own (content `filter` bleeds past the element box like CSS). + clip_rounded: f32, + // 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the origin + // so a stationary element blurs identically at every window size); 0.0 = 1:1 copy (scene blit). + downsample: f32, +} + +@group(1) @binding(0) var blur_locals: BlurParams; +@group(1) @binding(1) var t_blur: texture_2d; +@group(1) @binding(2) var s_blur: sampler; + +struct BlurVarying { + @builtin(position) position: vec4, + @location(0) uv: vec2, + @location(3) clip_distances: vec4, +} + +@vertex +fn vs_blur_fullscreen(@builtin(vertex_index) vertex_id: u32) -> BlurVarying { + // A single triangle large enough to cover the whole framebuffer. + let uv = vec2(f32((vertex_id << 1u) & 2u), f32(vertex_id & 2u)); + var out = BlurVarying(); + out.uv = uv; + out.position = vec4(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, 0.0, 1.0); + out.clip_distances = vec4(1.0); + return out; +} + +@fragment +fn fs_blur_downsample(input: BlurVarying) -> @location(0) vec4 { + if (blur_locals.downsample > 0.5) { + // Snapped 2:1 box downsample. Half-res texel `px` samples source at full-res coordinate + // 2*px + 1 (the boundary between source texels 2*px and 2*px+1), so one bilinear tap + // averages exactly that pair. The grid is anchored to the origin and independent of the + // viewport size, so an element at fixed pixels blurs identically at every window size — + // otherwise the implicit `floor(W/2)` grid stretches and the halo wobbles by ~1px on resize. + let dst = floor(input.position.xy); + let src_uv = (dst * 2.0 + 1.0) / globals.viewport_size; + return textureSampleLevel(t_blur, s_blur, src_uv, 0.0); + } + // 1:1 copy at matching resolution (used to blit the offscreen scene into the swapchain). + return textureSampleLevel(t_blur, s_blur, input.uv, 0.0); +} + +@fragment +fn fs_blur(input: BlurVarying) -> @location(0) vec4 { + let sigma = blur_locals.sigma; + let taps = i32(blur_locals.tap_count); + let step = blur_locals.tap_step; + var color = vec4(0.0); + var weight_sum = 0.0; + for (var i = -taps; i <= taps; i = i + 1) { + let offset = f32(i) * step; + let weight = gaussian(offset, sigma); + let uv = input.uv + blur_locals.direction * offset; + color += textureSampleLevel(t_blur, s_blur, uv, 0.0) * weight; + weight_sum += weight; + } + return color / max(weight_sum, 1e-5); +} + +@vertex +fn vs_blur_composite(@builtin(vertex_index) vertex_id: u32) -> BlurVarying { + let unit_vertex = vec2(f32(vertex_id & 1u), 0.5 * f32(vertex_id & 2u)); + var out = BlurVarying(); + out.position = to_device_position(unit_vertex, blur_locals.bounds); + out.uv = unit_vertex; + out.clip_distances = distance_from_clip_rect(unit_vertex, blur_locals.bounds, blur_locals.content_mask); + return out; +} + +@fragment +fn fs_blur_composite(input: BlurVarying) -> @location(0) vec4 { + if (any(input.clip_distances < vec4(0.0))) { + return vec4(0.0); + } + + // Sample the half-res blur by screen position, using the SAME fixed 2:1 grid the snapped + // downsample wrote (anchored at the origin, independent of viewport parity). `2*floor(W/2)` is + // the source span the half-res texture covers; dividing by it maps screen pixel p to half-res + // texel p/2 at every window size, so the composite stays put rather than wobbling on resize. + let blur_span = 2.0 * floor(globals.viewport_size * 0.5); + let uv = input.position.xy / blur_span; + let blurred = textureSampleLevel(t_blur, s_blur, uv, 0.0); + + let corner_radii = Corners( + blur_locals.corner_radii.x, + blur_locals.corner_radii.y, + blur_locals.corner_radii.z, + blur_locals.corner_radii.w, + ); + // Backdrop blur clips to the rounded rect (the frosted panel has a defined shape). Content + // blur does not — it bleeds past the element box like CSS `filter: blur`, so the soft fade + // isn't sharply truncated at the edge; its shape comes from the blurred group's own alpha. + let distance = quad_sdf(input.position.xy, blur_locals.bounds, corner_radii); + let coverage = select(1.0, saturate(0.5 - distance), blur_locals.clip_rounded > 0.5); + + // The blurred sample is premultiplied (blurring against the transparent, rgb=0 surround scales + // rgb with the fading alpha), so output premultiplied and let the pipeline blend premultiplied. + // A backdrop's scene is opaque (alpha ~= 1) so this replaces; a content-filter group is + // transparent outside its subtree, so the target shows through there instead of darkening. + let c = coverage * blur_locals.opacity; + return vec4(blurred.rgb * c, blurred.a * c); +} diff --git a/crates/gpui_wgpu/src/wgpu_renderer.rs b/crates/gpui_wgpu/src/wgpu_renderer.rs index 811fb39893..0b8ce0bede 100644 --- a/crates/gpui_wgpu/src/wgpu_renderer.rs +++ b/crates/gpui_wgpu/src/wgpu_renderer.rs @@ -1,11 +1,24 @@ use crate::{CompositorGpuHint, WgpuAtlas, WgpuContext, WgpuDeviceRequirements}; use bytemuck::{Pod, Zeroable}; use gpui::{ - AtlasTextureId, Background, Bounds, DevicePixels, GpuSpecs, MonochromeSprite, PaintSurface, - Path, Point, PolychromeSprite, PrimitiveBatch, Quad, ScaledPixels, Scene, Shadow, Size, - SubpixelSprite, Underline, get_gamma_correction_ratios, + AtlasTextureId, BackdropFilter, Background, Bounds, DevicePixels, FilterBoundary, GpuSpecs, + MonochromeSprite, PaintSurface, Path, Point, PolychromeSprite, PrimitiveBatch, Quad, + ScaledFilter, ScaledPixels, Scene, Shadow, Size, SubpixelSprite, Underline, + get_gamma_correction_ratios, }; use log::warn; + +/// The largest blur radius in a scene-space filter chain, in device pixels — used to size the +/// blur kernel and the dilated region the blur passes are scissored to. +/// +/// The `match` is exhaustive on purpose: adding a [`ScaledFilter`] variant breaks it here, +/// forcing this backend to handle (or deliberately ignore) the new filter rather than silently +/// dropping it. +fn max_blur_radius(filters: &[ScaledFilter]) -> f32 { + filters.iter().fold(0.0, |acc, filter| match filter { + ScaledFilter::Blur(radius) => acc.max(radius.0), + }) +} #[cfg(not(target_family = "wasm"))] use raw_window_handle::{HasDisplayHandle, HasWindowHandle}; use std::cell::RefCell; @@ -22,7 +35,7 @@ struct GlobalParams { } #[repr(C)] -#[derive(Clone, Copy, Pod, Zeroable)] +#[derive(Clone, Copy, Default, Pod, Zeroable)] struct PodBounds { origin: [f32; 2], size: [f32; 2], @@ -44,6 +57,38 @@ struct SurfaceParams { content_mask: PodBounds, } +/// Uniform passed to the blur pipelines. The same struct drives the downsample, separable +/// gaussian, and composite passes; fields not relevant to a given pass are left zero. +#[repr(C)] +#[derive(Clone, Copy, Default, Pod, Zeroable)] +struct BlurParams { + /// Composite target rectangle, in device pixels (composite pass only). + bounds: PodBounds, + /// Clip rectangle, in device pixels (composite pass only). + content_mask: PodBounds, + /// Rounded-corner radii (tl, tr, br, bl), in device pixels (composite pass only). + corner_radii: [f32; 4], + /// Per-tap sampling step in UV space (gaussian passes only): (1/width, 0) or (0, 1/height). + direction: [f32; 2], + /// Gaussian sigma, in the (half-resolution) blur texture's pixels. + sigma: f32, + /// Element opacity, multiplied into the composited result. + opacity: f32, + /// Number of taps to each side of center (gaussian passes only). + tap_count: f32, + /// Spacing between taps in pixels; >1 lets `tap_count` taps span very large radii without + /// truncating the gaussian (see #6 in review). + tap_step: f32, + /// 1.0 to clip the composite to the rounded rect (backdrop — the panel has a defined shape), + /// 0.0 to let the blurred result fade out on its own (content `filter` — it bleeds past the + /// element bounds like CSS, so the fade isn't sharply truncated at the box edge). + clip_rounded: f32, + /// 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the + /// origin, so a stationary element blurs identically at every window size); 0.0 = 1:1 copy + /// (the scene blit, which must not downsample). Downsample pass only. + downsample: f32, +} + #[repr(C)] #[derive(Clone, Copy, Pod, Zeroable)] struct GammaParams { @@ -91,6 +136,14 @@ struct WgpuPipelines { subpixel_sprites: Option, poly_sprites: wgpu::RenderPipeline, surfaces: wgpu::RenderPipeline, + /// Copies a source texture into the (smaller) target with one bilinear tap. Used both to + /// downsample the scene into the half-resolution blur texture and to blit the offscreen + /// scene into the swapchain at the end of the frame. + blur_downsample: wgpu::RenderPipeline, + /// One axis of a separable gaussian blur; direction is supplied per draw via [`BlurParams`]. + blur: wgpu::RenderPipeline, + /// Composites a blurred texture into a rounded rectangle (with clip + opacity). + blur_composite: wgpu::RenderPipeline, } struct WgpuBindGroupLayouts { @@ -98,6 +151,7 @@ struct WgpuBindGroupLayouts { instances: wgpu::BindGroupLayout, instances_with_texture: wgpu::BindGroupLayout, surfaces: wgpu::BindGroupLayout, + blur: wgpu::BindGroupLayout, } /// Shared GPU context reference, used to coordinate device recovery across multiple windows. @@ -113,6 +167,10 @@ struct WgpuResources { atlas_sampler: wgpu::Sampler, surface_sampler: wgpu::Sampler, surface_uniform_buffer: wgpu::Buffer, + /// One reused uniform buffer holding [`BlurParams`] for every blur pass in a frame, each at a + /// distinct (alignment-strided) offset. Avoids allocating a buffer per pass; distinct offsets + /// mean `write_buffer`'s last-write-at-submit semantics don't clobber earlier passes. + blur_params_buffer: wgpu::Buffer, globals_buffer: wgpu::Buffer, globals_bind_group: wgpu::BindGroup, path_globals_bind_group: wgpu::BindGroup, @@ -121,6 +179,23 @@ struct WgpuResources { path_intermediate_view: Option, path_msaa_texture: Option, path_msaa_view: Option, + /// Blur offscreen targets. Allocated lazily (only when a frame actually uses a blur filter) + /// so apps that never blur pay no extra VRAM. `None`/empty until first use. + /// + /// Full-resolution offscreen color target the scene is rendered into so that blur passes + /// can sample already-painted content; blitted to the swapchain at the end of the frame. + scene_color_texture: Option, + scene_color_view: Option, + /// Half-resolution ping/pong targets for the downsample + separable gaussian passes. + blur_ping_texture: Option, + blur_ping_view: Option, + blur_pong_texture: Option, + blur_pong_view: Option, + /// Full-resolution offscreen targets a content-filter (`filter`) group renders into before + /// being blurred and composited back. One per nesting level (indexed by depth) so nested + /// content blurs isolate correctly, up to [`MAX_FILTER_DEPTH`]; deeper nests render inline. + group_textures: Vec, + group_views: Vec, } impl WgpuResources { @@ -129,9 +204,26 @@ impl WgpuResources { self.path_intermediate_view = None; self.path_msaa_texture = None; self.path_msaa_view = None; + self.scene_color_texture = None; + self.scene_color_view = None; + self.blur_ping_texture = None; + self.blur_ping_view = None; + self.blur_pong_texture = None; + self.blur_pong_view = None; + self.group_textures.clear(); + self.group_views.clear(); } } +/// Number of content-filter (`filter`) nesting levels that get their own isolated group texture. +/// Two covers the realistic "a blurred element inside another blurred element" case; deeper nests +/// render inline (unblurred at the inner level) rather than allocating unbounded VRAM. +const MAX_FILTER_DEPTH: usize = 2; + +/// Number of [`BlurParams`] slots in the shared blur-params buffer (one per blur pass per frame). +/// Each frame uses 4 passes per backdrop/group plus one blit; 256 covers dozens of filters. +const BLUR_PARAMS_SLOTS: u64 = 256; + pub struct WgpuRenderer { /// Shared GPU context for device recovery coordination (unused on WASM). #[allow(dead_code)] @@ -150,6 +242,10 @@ pub struct WgpuRenderer { instance_buffer_capacity: u64, max_buffer_size: u64, storage_buffer_alignment: u64, + /// Stride between [`BlurParams`] slots in `blur_params_buffer`, and a per-frame bump cursor + /// (in slots) handed out to blur passes. Cell so the `&self` blur helpers can advance it. + blur_params_stride: u64, + blur_params_slot: std::cell::Cell, rendering_params: RenderingParameters, is_bgr: bool, dual_source_blending: bool, @@ -395,6 +491,16 @@ impl WgpuRenderer { }); let uniform_alignment = device.limits().min_uniform_buffer_offset_alignment as u64; + // Shared blur-params buffer: BLUR_PARAMS_SLOTS slots, each one alignment stride apart. + let blur_params_stride = + (std::mem::size_of::() as u64).next_multiple_of(uniform_alignment); + let blur_params_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("blur_params_buffer"), + size: blur_params_stride * BLUR_PARAMS_SLOTS, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let globals_size = std::mem::size_of::() as u64; let gamma_size = std::mem::size_of::() as u64; let path_globals_offset = globals_size.next_multiple_of(uniform_alignment); @@ -481,6 +587,7 @@ impl WgpuRenderer { atlas_sampler, surface_sampler, surface_uniform_buffer, + blur_params_buffer, globals_buffer, globals_bind_group, path_globals_bind_group, @@ -491,6 +598,14 @@ impl WgpuRenderer { path_intermediate_view: None, path_msaa_texture: None, path_msaa_view: None, + scene_color_texture: None, + scene_color_view: None, + blur_ping_texture: None, + blur_ping_view: None, + blur_pong_texture: None, + blur_pong_view: None, + group_textures: Vec::new(), + group_views: Vec::new(), }; Ok(Self { @@ -505,6 +620,8 @@ impl WgpuRenderer { instance_buffer_capacity: initial_instance_buffer_capacity, max_buffer_size, storage_buffer_alignment, + blur_params_stride, + blur_params_slot: std::cell::Cell::new(0), rendering_params, is_bgr: false, dual_source_blending, @@ -626,11 +743,44 @@ impl WgpuRenderer { ], }); + let blur = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("blur_layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::VERTEX_FRAGMENT, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: NonZeroU64::new(std::mem::size_of::() as u64), + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + ], + }); + WgpuBindGroupLayouts { globals, instances, instances_with_texture, surfaces, + blur, } } @@ -891,7 +1041,60 @@ impl WgpuRenderer { &layouts.globals, &layouts.surfaces, wgpu::PrimitiveTopology::TriangleStrip, - &[Some(color_target)], + &[Some(color_target.clone())], + 1, + &shader_module, + ); + + // Blur pipelines all sample one texture into another; the downsample and gaussian passes + // overwrite their (intermediate) target, while the composite blends over the scene. + let no_blend_target = wgpu::ColorTargetState { + format: surface_format, + blend: None, + write_mask: wgpu::ColorWrites::ALL, + }; + + let blur_downsample = create_pipeline( + "blur_downsample", + "vs_blur_fullscreen", + "fs_blur_downsample", + &layouts.globals, + &layouts.blur, + wgpu::PrimitiveTopology::TriangleList, + &[Some(no_blend_target.clone())], + 1, + &shader_module, + ); + + let blur = create_pipeline( + "blur", + "vs_blur_fullscreen", + "fs_blur", + &layouts.globals, + &layouts.blur, + wgpu::PrimitiveTopology::TriangleList, + &[Some(no_blend_target)], + 1, + &shader_module, + ); + + // The blurred sample is premultiplied (blurring against the transparent, rgb=0 region + // around the source scales rgb with the fading alpha), so the composite outputs + // premultiplied and blends premultiplied — straight alpha blending would multiply rgb by + // alpha a second time and darken the faded edges. Independent of the window's alpha mode. + let premultiplied_target = wgpu::ColorTargetState { + format: surface_format, + blend: Some(wgpu::BlendState::PREMULTIPLIED_ALPHA_BLENDING), + write_mask: wgpu::ColorWrites::ALL, + }; + let blur_composite = create_pipeline( + "blur_composite", + "vs_blur_composite", + "fs_blur_composite", + &layouts.globals, + &layouts.blur, + wgpu::PrimitiveTopology::TriangleStrip, + &[Some(premultiplied_target)], 1, &shader_module, ); @@ -906,6 +1109,9 @@ impl WgpuRenderer { subpixel_sprites, poly_sprites, surfaces, + blur_downsample, + blur, + blur_composite, } } @@ -998,6 +1204,19 @@ impl WgpuRenderer { if let Some(ref texture) = resources.path_msaa_texture { texture.destroy(); } + for texture in [ + &resources.scene_color_texture, + &resources.blur_ping_texture, + &resources.blur_pong_texture, + ] + .into_iter() + .flatten() + { + texture.destroy(); + } + for texture in &resources.group_textures { + texture.destroy(); + } resources .surface @@ -1038,6 +1257,40 @@ impl WgpuRenderer { resources.path_msaa_view = path_msaa_view; } + /// Lazily allocate the blur offscreen targets — the full-res scene texture, half-res + /// ping/pong, and one full-res group texture per nesting level. Called only on frames that + /// actually use a blur filter, so non-blurring apps never pay this VRAM. A no-op once + /// allocated (invalidated alongside the path intermediates on resize / device loss). + fn ensure_blur_textures(&mut self) { + if self.resources().scene_color_texture.is_some() { + return; + } + let format = self.surface_config.format; + let width = self.surface_config.width; + let height = self.surface_config.height; + let blur_width = (width / 2).max(1); + let blur_height = (height / 2).max(1); + let resources = self.resources_mut(); + + let (t, v) = Self::create_path_intermediate(&resources.device, format, width, height); + resources.scene_color_texture = Some(t); + resources.scene_color_view = Some(v); + let (t, v) = + Self::create_path_intermediate(&resources.device, format, blur_width, blur_height); + resources.blur_ping_texture = Some(t); + resources.blur_ping_view = Some(v); + let (t, v) = + Self::create_path_intermediate(&resources.device, format, blur_width, blur_height); + resources.blur_pong_texture = Some(t); + resources.blur_pong_view = Some(v); + + for _ in 0..MAX_FILTER_DEPTH { + let (t, v) = Self::create_path_intermediate(&resources.device, format, width, height); + resources.group_textures.push(t); + resources.group_views.push(v); + } + } + pub fn set_subpixel_layout(&mut self, is_bgr: bool) { self.is_bgr = is_bgr; } @@ -1171,6 +1424,14 @@ impl WgpuRenderer { // Now that we know the surface is healthy, ensure intermediate textures exist self.ensure_intermediate_textures(); + // Blur is the only thing that needs the offscreen scene texture; allocate it (and the + // ping/pong/group targets) lazily so non-blurring apps pay no extra VRAM or blit. + let use_offscreen = + !scene.backdrop_filters.is_empty() || !scene.filter_boundaries.is_empty(); + if use_offscreen { + self.ensure_blur_textures(); + } + let frame_view = frame .texture .create_view(&wgpu::TextureViewDescriptor::default()); @@ -1224,6 +1485,8 @@ impl WgpuRenderer { loop { let mut instance_offset: u64 = 0; + // Reset the blur-params bump cursor each (re)render of the scene. + self.blur_params_slot.set(0); let mut overflow = false; let mut encoder = @@ -1233,11 +1496,41 @@ impl WgpuRenderer { label: Some("main_encoder"), }); + // When the scene contains blur filters, render into the offscreen scene texture (so + // filters can sample already-painted content mid-frame) and blit to the swapchain at + // the end; otherwise render straight to the swapchain. `use_offscreen` and the blur + // textures were computed/allocated above. + let scene_color_view = if use_offscreen { + Some( + self.resources() + .scene_color_view + .as_ref() + .expect("scene_color_view allocated by ensure_blur_textures") + .clone(), + ) + } else { + None + }; + // The active render target. While inside a content-filter (`filter`) group it points + // at a group texture so the group renders in isolation. + let mut current_target = match &scene_color_view { + Some(view) => view.clone(), + None => frame_view.clone(), + }; + // One group texture per nesting depth; empty when not blurring. + let group_views = if use_offscreen { + self.resources().group_views.clone() + } else { + Vec::new() + }; + // (boundary, parent target to composite back into, whether this level is isolated). + let mut filter_stack: Vec<(FilterBoundary, wgpu::TextureView, bool)> = Vec::new(); + { let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("main_pass"), color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &frame_view, + view: ¤t_target, resolve_target: None, ops: wgpu::Operations { load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), @@ -1276,7 +1569,7 @@ impl WgpuRenderer { pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("main_pass_continued"), color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &frame_view, + view: ¤t_target, resolve_target: None, ops: wgpu::Operations { load: wgpu::LoadOp::Load, @@ -1327,6 +1620,103 @@ impl WgpuRenderer { PrimitiveBatch::Surfaces(range) => { self.draw_surfaces(&scene.surfaces[range], &mut pass) } + PrimitiveBatch::BackdropFilters(range) => { + // Interrupt the current pass, blur the content painted so far behind + // each backdrop's rounded rect, then resume drawing on top. + drop(pass); + for filter in &scene.backdrop_filters[range] { + self.draw_backdrop_filter(&mut encoder, filter, ¤t_target); + } + pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("main_pass_continued"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: ¤t_target, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); + true + } + PrimitiveBatch::FilterBoundary(ix) => { + let boundary = scene.filter_boundaries[ix].clone(); + if boundary.is_start { + // Each isolated nesting level uses its own group texture from the + // pool (indexed by current isolation depth). Beyond the pool size + // (MAX_FILTER_DEPTH) deeper filters render inline without isolation + // rather than corrupting an outer group. + let depth = filter_stack.iter().filter(|entry| entry.2).count(); + if depth < group_views.len() { + drop(pass); + let parent = current_target.clone(); + current_target = group_views[depth].clone(); + filter_stack.push((boundary, parent, true)); + pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("filter_group"), + color_attachments: &[Some( + wgpu::RenderPassColorAttachment { + view: ¤t_target, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear( + wgpu::Color::TRANSPARENT, + ), + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + }, + )], + depth_stencil_attachment: None, + ..Default::default() + }); + } else { + filter_stack.push((boundary, current_target.clone(), false)); + } + } else if let Some((boundary, parent, isolated)) = filter_stack.pop() { + if isolated { + drop(pass); + self.blur_and_composite( + &mut encoder, + ¤t_target, + &parent, + boundary.bounds, + boundary.content_mask.bounds, + [ + boundary.corner_radii.top_left.0, + boundary.corner_radii.top_right.0, + boundary.corner_radii.bottom_right.0, + boundary.corner_radii.bottom_left.0, + ], + max_blur_radius(&boundary.filters), + boundary.opacity, + false, + ); + current_target = parent; + pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("main_pass_continued"), + color_attachments: &[Some( + wgpu::RenderPassColorAttachment { + view: ¤t_target, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + }, + )], + depth_stencil_attachment: None, + ..Default::default() + }); + } + } + true + } }; if !ok { overflow = true; @@ -1349,6 +1739,12 @@ impl WgpuRenderer { continue; } + // Present the offscreen scene by copying it into the swapchain texture. Skipped when + // rendering went straight to the swapchain (no filters this frame). + if let Some(scene_color_view) = &scene_color_view { + self.blit_to_frame(&mut encoder, scene_color_view, &frame_view); + } + self.resources() .queue .submit(std::iter::once(encoder.finish())); @@ -1498,6 +1894,289 @@ impl WgpuRenderer { true } + /// Build a bind group for a blur pass. Writes `params` into the next slot of the shared + /// `blur_params_buffer` (no per-pass allocation) and references that slot, the source texture, + /// and the filtering sampler. Distinct per-pass offsets keep `write_buffer`'s + /// last-write-at-submit semantics from clobbering earlier passes within a frame. + fn make_blur_bind_group( + &self, + params: BlurParams, + source: &wgpu::TextureView, + ) -> wgpu::BindGroup { + let resources = self.resources(); + let slot = self.blur_params_slot.get() % BLUR_PARAMS_SLOTS; + self.blur_params_slot.set(slot + 1); + let offset = slot * self.blur_params_stride; + resources.queue.write_buffer( + &resources.blur_params_buffer, + offset, + bytemuck::bytes_of(¶ms), + ); + resources + .device + .create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("blur_bind_group"), + layout: &resources.bind_group_layouts.blur, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { + buffer: &resources.blur_params_buffer, + offset, + size: NonZeroU64::new(std::mem::size_of::() as u64), + }), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::TextureView(source), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: wgpu::BindingResource::Sampler(&resources.surface_sampler), + }, + ], + }) + } + + /// Run a full-screen (3-vertex) blur pass that overwrites `target` by sampling `source`. + /// `scissor` (x, y, w, h, in `target` pixels) limits fragment work to the region that + /// actually feeds the composite — the element bounds dilated by the kernel radius. + fn run_blur_pass( + &self, + encoder: &mut wgpu::CommandEncoder, + label: &str, + pipeline: &wgpu::RenderPipeline, + target: &wgpu::TextureView, + source: &wgpu::TextureView, + params: BlurParams, + scissor: [u32; 4], + ) { + let bind_group = self.make_blur_bind_group(params, source); + let resources = self.resources(); + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some(label), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: target, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); + pass.set_pipeline(pipeline); + pass.set_bind_group(0, &resources.globals_bind_group, &[]); + pass.set_bind_group(1, &bind_group, &[]); + pass.set_scissor_rect(scissor[0], scissor[1], scissor[2], scissor[3]); + pass.draw(0..3, 0..1); + } + + /// Blur `source` (full-resolution) and composite the result into `target`, clipped to + /// `bounds`/`corner_radii`/`content_mask` and modulated by `opacity`. Shared by the backdrop + /// and content-filter paths. Uses the half-resolution ping/pong textures as scratch. + #[allow(clippy::too_many_arguments)] + fn blur_and_composite( + &self, + encoder: &mut wgpu::CommandEncoder, + source: &wgpu::TextureView, + target: &wgpu::TextureView, + bounds: Bounds, + content_mask: Bounds, + corner_radii: [f32; 4], + blur_radius: f32, + opacity: f32, + // Backdrop clips to the rounded rect; content (`filter`) bleeds past its bounds. + clip_rounded: bool, + ) { + // Sigma is halved because the blur runs at half resolution. + let sigma = (blur_radius * 0.5).max(0.0); + if sigma <= 0.0 { + return; + } + // Span ±3σ. If that needs more than 32 taps, spread the taps apart (tap_step > 1) rather + // than truncating the kernel — keeps very large radii from clipping (review #6). + let ideal_taps = (3.0 * sigma).ceil(); + let tap_count = ideal_taps.clamp(1.0, 32.0); + let tap_step = (ideal_taps / tap_count).max(1.0); + let full_w = self.surface_config.width; + let full_h = self.surface_config.height; + let blur_width = (full_w / 2).max(1) as f32; + let blur_height = (full_h / 2).max(1) as f32; + + // Limit the half-res passes to the element bounds dilated by the kernel radius (3·sigma, + // full-res) — outside that the composite never samples, so there's no reason to blur it. + let dilation = 3.0 * blur_radius; + let hw = (full_w / 2).max(1); + let hh = (full_h / 2).max(1); + let x0 = (((bounds.origin.x.0 - dilation) * 0.5).floor().max(0.0) as u32).min(hw); + let y0 = (((bounds.origin.y.0 - dilation) * 0.5).floor().max(0.0) as u32).min(hh); + let x1 = ((((bounds.origin.x.0 + bounds.size.width.0 + dilation) * 0.5) + .ceil() + .max(0.0) as u32) + .min(hw)) + .max(x0); + let y1 = ((((bounds.origin.y.0 + bounds.size.height.0 + dilation) * 0.5) + .ceil() + .max(0.0) as u32) + .min(hh)) + .max(y0); + let scissor = [x0, y0, x1 - x0, y1 - y0]; + if scissor[2] == 0 || scissor[3] == 0 { + return; + } + + // Owned handles so the passes below don't borrow `self`. + let (ping, pong) = { + let resources = self.resources(); + match ( + resources.blur_ping_view.as_ref(), + resources.blur_pong_view.as_ref(), + ) { + (Some(ping), Some(pong)) => (ping.clone(), pong.clone()), + _ => return, + } + }; + + // Downsample source -> ping, then separable gaussian ping -> pong -> ping. + self.run_blur_pass( + encoder, + "blur_downsample", + &self.resources().pipelines.blur_downsample, + &ping, + source, + BlurParams { + downsample: 1.0, + ..Default::default() + }, + scissor, + ); + self.run_blur_pass( + encoder, + "blur_horizontal", + &self.resources().pipelines.blur, + &pong, + &ping, + BlurParams { + direction: [1.0 / blur_width, 0.0], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + scissor, + ); + self.run_blur_pass( + encoder, + "blur_vertical", + &self.resources().pipelines.blur, + &ping, + &pong, + BlurParams { + direction: [0.0, 1.0 / blur_height], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + scissor, + ); + + // Composite the blurred result into the target (loads existing content). For content blur + // the quad covers the dilated region so the blur can fade out past the element box (no + // sharp clip); for backdrop the quad is the element bounds and the shader clips to the + // rounded rect. + let composite_bounds = if clip_rounded { + bounds + } else { + bounds.dilate(ScaledPixels(dilation)) + }; + let params = BlurParams { + bounds: composite_bounds.into(), + content_mask: content_mask.into(), + corner_radii, + opacity, + clip_rounded: if clip_rounded { 1.0 } else { 0.0 }, + ..Default::default() + }; + let bind_group = self.make_blur_bind_group(params, &ping); + let resources = self.resources(); + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("blur_composite"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: target, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); + pass.set_pipeline(&resources.pipelines.blur_composite); + pass.set_bind_group(0, &resources.globals_bind_group, &[]); + pass.set_bind_group(1, &bind_group, &[]); + pass.draw(0..4, 0..1); + } + + /// Blur the scene painted so far behind `filter.bounds` and composite it back as frosted glass. + fn draw_backdrop_filter( + &self, + encoder: &mut wgpu::CommandEncoder, + filter: &BackdropFilter, + scene_color_view: &wgpu::TextureView, + ) { + self.blur_and_composite( + encoder, + scene_color_view, + scene_color_view, + filter.bounds, + filter.content_mask.bounds, + [ + filter.corner_radii.top_left.0, + filter.corner_radii.top_right.0, + filter.corner_radii.bottom_right.0, + filter.corner_radii.bottom_left.0, + ], + max_blur_radius(&filter.filters), + filter.opacity, + true, + ); + } + + /// Copy the offscreen scene texture into the swapchain texture. + fn blit_to_frame( + &self, + encoder: &mut wgpu::CommandEncoder, + source: &wgpu::TextureView, + frame_view: &wgpu::TextureView, + ) { + let bind_group = self.make_blur_bind_group(BlurParams::default(), source); + let resources = self.resources(); + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("scene_blit"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: frame_view, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); + pass.set_pipeline(&resources.pipelines.blur_downsample); + pass.set_bind_group(0, &resources.globals_bind_group, &[]); + pass.set_bind_group(1, &bind_group, &[]); + pass.draw(0..3, 0..1); + } + fn draw_polychrome_sprites( &self, sprites: &[PolychromeSprite], diff --git a/crates/gpui_windows/build.rs b/crates/gpui_windows/build.rs index 1db95715e2..4d93b0a417 100644 --- a/crates/gpui_windows/build.rs +++ b/crates/gpui_windows/build.rs @@ -38,6 +38,9 @@ mod shader_compilation { "monochrome_sprite", "subpixel_sprite", "polychrome_sprite", + "blur_downsample", + "blur", + "blur_composite", ]; let rust_binding_path = format!("{}/shaders_bytes.rs", out_dir); diff --git a/crates/gpui_windows/src/directx_renderer.rs b/crates/gpui_windows/src/directx_renderer.rs index b5c4d3bf34..61ec8355fc 100644 --- a/crates/gpui_windows/src/directx_renderer.rs +++ b/crates/gpui_windows/src/directx_renderer.rs @@ -23,11 +23,29 @@ use crate::directx_renderer::shader_resources::{RawShaderBytes, ShaderModule, Sh use crate::*; use gpui::*; +/// The largest blur radius in a scene-space filter chain, in device pixels — used to size the +/// blur kernel and the dilated region the blur passes are scissored to. +/// +/// The `match` is exhaustive on purpose: adding a [`ScaledFilter`] variant breaks it here, +/// forcing this backend to handle (or deliberately ignore) the new filter rather than silently +/// dropping it. +fn max_blur_radius(filters: &[ScaledFilter]) -> f32 { + filters.iter().fold(0.0, |acc, filter| match filter { + ScaledFilter::Blur(radius) => acc.max(radius.0), + }) +} + pub(crate) const DISABLE_DIRECT_COMPOSITION: &str = "GPUI_DISABLE_DIRECT_COMPOSITION"; const RENDER_TARGET_FORMAT: DXGI_FORMAT = DXGI_FORMAT_B8G8R8A8_UNORM; // This configuration is used for MSAA rendering on paths only, and it's guaranteed to be supported by DirectX 11. const PATH_MULTISAMPLE_COUNT: u32 = 4; +/// Number of content-filter (`filter`) nesting levels that get their own isolated group target. +/// Two covers the realistic "a blurred element inside another blurred element" case; deeper nests +/// render inline (unblurred at the inner level) rather than allocating unbounded VRAM. Must match +/// the wgpu backend's `MAX_FILTER_DEPTH` so nested blur renders consistently across platforms. +const MAX_FILTER_DEPTH: usize = 2; + pub(crate) struct FontInfo { pub gamma_ratios: [f32; 4], pub grayscale_enhanced_contrast: f32, @@ -53,6 +71,12 @@ pub(crate) struct DirectXRenderer { /// In that case we want to discard the first frame that we draw as we got reset in the middle of a frame /// meaning we lost all the allocated gpu textures and scene resources. skip_draws: bool, + + /// The render target currently bound for the main scene this frame (the offscreen + /// `scene_color` when blur filters are present, a content-filter group texture inside such a + /// group, or the swapchain otherwise). `draw_paths_to_intermediate` restores to this after + /// its own pass so paths land on the correct target. + active_render_target: Option, } /// Direct3D objects @@ -77,10 +101,68 @@ struct DirectXResources { path_intermediate_msaa_texture: ID3D11Texture2D, path_intermediate_msaa_view: Option, + // Offscreen targets for blur filters (each is render-target + shader-resource). + blur: BlurResources, + // Cached viewport viewport: D3D11_VIEWPORT, } +/// Offscreen render targets used by the blur filters. The scene is rendered into `scene_color` +/// (so filters can sample it), `ping`/`pong` are half-resolution scratch for the separable +/// gaussian, and `groups` isolate content-filter (`filter`) subtrees — one per nesting level +/// (indexed by isolation depth), up to [`MAX_FILTER_DEPTH`], so nested content blurs isolate +/// correctly; deeper nests render inline. +struct BlurResources { + scene_color: ID3D11Texture2D, + scene_color_rtv: Option, + scene_color_srv: Option, + ping: ID3D11Texture2D, + ping_rtv: Option, + ping_srv: Option, + pong: ID3D11Texture2D, + pong_rtv: Option, + pong_srv: Option, + // Kept alive for the lifetime of their views; indexed by isolation depth. + groups: Vec, + group_rtvs: Vec>, + group_srvs: Vec>, +} + +impl BlurResources { + fn new(device: &ID3D11Device, width: u32, height: u32) -> Result { + let half_w = (width / 2).max(1); + let half_h = (height / 2).max(1); + let (scene_color, scene_color_rtv, scene_color_srv) = + create_color_target(device, width, height)?; + let (ping, ping_rtv, ping_srv) = create_color_target(device, half_w, half_h)?; + let (pong, pong_rtv, pong_srv) = create_color_target(device, half_w, half_h)?; + let mut groups = Vec::with_capacity(MAX_FILTER_DEPTH); + let mut group_rtvs = Vec::with_capacity(MAX_FILTER_DEPTH); + let mut group_srvs = Vec::with_capacity(MAX_FILTER_DEPTH); + for _ in 0..MAX_FILTER_DEPTH { + let (group, group_rtv, group_srv) = create_color_target(device, width, height)?; + groups.push(group); + group_rtvs.push(group_rtv); + group_srvs.push(group_srv); + } + Ok(Self { + scene_color, + scene_color_rtv, + scene_color_srv, + ping, + ping_rtv, + ping_srv, + pong, + pong_rtv, + pong_srv, + groups, + group_rtvs, + group_srvs, + }) + } +} + struct DirectXRenderPipelines { shadow_pipeline: PipelineState, quad_pipeline: PipelineState, @@ -90,6 +172,18 @@ struct DirectXRenderPipelines { mono_sprites: PipelineState, subpixel_sprites: PipelineState, poly_sprites: PipelineState, + // Blur (backdrop-filter / filter). These don't use the generic PipelineState since they + // sample a texture rather than read a structured instance buffer; their parameters live in + // a dedicated constant buffer at register b1. + blur_downsample_vertex: ID3D11VertexShader, + blur_downsample_fragment: ID3D11PixelShader, + blur_vertex: ID3D11VertexShader, + blur_fragment: ID3D11PixelShader, + blur_composite_vertex: ID3D11VertexShader, + blur_composite_fragment: ID3D11PixelShader, + blur_params_buffer: ID3D11Buffer, + blur_blend_replace: ID3D11BlendState, + blur_blend_composite: ID3D11BlendState, } struct DirectXGlobalElements { @@ -174,6 +268,7 @@ impl DirectXRenderer { width: 1, height: 1, skip_draws: false, + active_render_target: None, }) } @@ -318,6 +413,56 @@ impl DirectXRenderer { self.upload_scene_buffers(scene)?; + // Only route through the offscreen scene texture when the scene contains blur filters; + // otherwise render straight to the swapchain exactly as before. + let use_offscreen = + !scene.backdrop_filters.is_empty() || !scene.filter_boundaries.is_empty(); + + // Clone the views we need (AddRef) so the loop can rebind render targets without holding a + // borrow of `self` across the `&mut self` draw_* calls. + let (scene_rtv, scene_srv, group_rtvs, group_srvs, swapchain_rtv) = { + let r = self.resources.as_ref().context("resources missing")?; + ( + r.blur.scene_color_rtv.clone(), + r.blur.scene_color_srv.clone(), + r.blur.group_rtvs.clone(), + r.blur.group_srvs.clone(), + r.render_target_view.clone(), + ) + }; + let ctx = self + .devices + .as_ref() + .context("devices missing")? + .device_context + .clone(); + + if use_offscreen { + unsafe { + if let Some(rtv) = scene_rtv.as_ref() { + ctx.ClearRenderTargetView(rtv, &[0.0; 4]); + } + ctx.OMSetRenderTargets(Some(slice::from_ref(&scene_rtv)), None); + } + self.active_render_target = scene_rtv.clone(); + } else { + self.active_render_target = swapchain_rtv.clone(); + } + + // Current target for the main scene + a parent stack for content-filter groups. + let mut current_rtv = self.active_render_target.clone(); + let mut current_srv = if use_offscreen { + scene_srv.clone() + } else { + None + }; + // (parent_rtv, parent_srv, isolated) + let mut filter_stack: Vec<( + Option, + Option, + bool, + )> = Vec::new(); + for batch in scene.batches() { match batch { PrimitiveBatch::Shadows(range) => self.draw_shadows(range.start, range.len()), @@ -338,6 +483,77 @@ impl DirectXRenderer { self.draw_polychrome_sprites(texture_id, range.start, range.len()) } PrimitiveBatch::Surfaces(range) => self.draw_surfaces(&scene.surfaces[range]), + PrimitiveBatch::BackdropFilters(range) => { + let result = (|| { + for filter in &scene.backdrop_filters[range] { + self.dx_blur_and_composite( + ¤t_srv, + ¤t_rtv, + filter.bounds, + filter.content_mask.bounds, + corner_radii_array(filter.corner_radii), + max_blur_radius(&filter.filters), + filter.opacity, + true, + )?; + } + Ok::<(), anyhow::Error>(()) + })(); + // Restore the current target for subsequent batches. + unsafe { + ctx.OMSetRenderTargets(Some(slice::from_ref(¤t_rtv)), None); + } + result + } + PrimitiveBatch::FilterBoundary(ix) => { + let boundary = scene.filter_boundaries[ix].clone(); + if boundary.is_start { + // Each isolated nesting level uses its own group target from the pool + // (indexed by current isolation depth). Beyond the pool size + // (MAX_FILTER_DEPTH) deeper filters render inline without isolation rather + // than corrupting an outer group. + let depth = filter_stack.iter().filter(|entry| entry.2).count(); + if depth < group_rtvs.len() { + filter_stack.push((current_rtv.clone(), current_srv.clone(), true)); + current_rtv = group_rtvs[depth].clone(); + current_srv = group_srvs[depth].clone(); + self.active_render_target = current_rtv.clone(); + unsafe { + if let Some(rtv) = current_rtv.as_ref() { + ctx.ClearRenderTargetView(rtv, &[0.0; 4]); + } + ctx.OMSetRenderTargets(Some(slice::from_ref(¤t_rtv)), None); + } + } else { + filter_stack.push((current_rtv.clone(), current_srv.clone(), false)); + } + Ok(()) + } else if let Some((parent_rtv, parent_srv, isolated)) = filter_stack.pop() { + let result = if isolated { + self.dx_blur_and_composite( + ¤t_srv, + &parent_rtv, + boundary.bounds, + boundary.content_mask.bounds, + corner_radii_array(boundary.corner_radii), + max_blur_radius(&boundary.filters), + boundary.opacity, + false, + ) + } else { + Ok(()) + }; + current_rtv = parent_rtv; + current_srv = parent_srv; + self.active_render_target = current_rtv.clone(); + unsafe { + ctx.OMSetRenderTargets(Some(slice::from_ref(¤t_rtv)), None); + } + result + } else { + Ok(()) + } + } } .context(format!( "scene too large:\ @@ -352,6 +568,12 @@ impl DirectXRenderer { scene.surfaces.len(), ))?; } + + // Present the offscreen scene by blitting it into the swapchain. + if use_offscreen { + self.dx_blit(&scene_srv, &swapchain_rtv)?; + } + self.active_render_target = None; self.present() } @@ -553,10 +775,16 @@ impl DirectXRenderer { 0, RENDER_TARGET_FORMAT, ); - // Restore main render target + // Restore the active render target (the offscreen scene/group target when blurring, + // otherwise the swapchain) so the path sprites land on the correct surface. + let restore_target = if self.active_render_target.is_some() { + &self.active_render_target + } else { + &resources.render_target_view + }; devices .device_context - .OMSetRenderTargets(Some(slice::from_ref(&resources.render_target_view)), None); + .OMSetRenderTargets(Some(slice::from_ref(restore_target)), None); } Ok(()) @@ -704,6 +932,208 @@ impl DirectXRenderer { Ok(()) } + /// Run a single blur pass: a full-screen (or composite) draw sampling `source_srv` into + /// `target_rtv`, with `params` in the blur constant buffer (b1). + #[allow(clippy::too_many_arguments)] + fn dx_blur_pass( + &self, + vertex: &ID3D11VertexShader, + fragment: &ID3D11PixelShader, + blend: &ID3D11BlendState, + target_rtv: &Option, + source_srv: &Option, + params: BlurParams, + viewport: &D3D11_VIEWPORT, + topology: D3D_PRIMITIVE_TOPOLOGY, + vertex_count: u32, + clear: bool, + ) -> Result<()> { + let devices = self.devices.as_ref().context("devices missing")?; + let ctx = &devices.device_context; + update_buffer(ctx, &self.pipelines.blur_params_buffer, &[params])?; + let null_srv: [Option; 1] = [None]; + let blur_params = [Some(self.pipelines.blur_params_buffer.clone())]; + unsafe { + // Unbind any SRV at slot 0 so the target texture isn't simultaneously bound as input. + ctx.PSSetShaderResources(0, Some(&null_srv)); + if clear { + ctx.ClearRenderTargetView( + target_rtv.as_ref().context("blur target view missing")?, + &[0.0; 4], + ); + } + ctx.OMSetRenderTargets(Some(slice::from_ref(target_rtv)), None); + ctx.RSSetViewports(Some(slice::from_ref(viewport))); + ctx.IASetPrimitiveTopology(topology); + ctx.VSSetShader(vertex, None); + ctx.PSSetShader(fragment, None); + ctx.VSSetConstantBuffers(0, Some(slice::from_ref(&self.globals.global_params_buffer))); + ctx.PSSetConstantBuffers(0, Some(slice::from_ref(&self.globals.global_params_buffer))); + ctx.VSSetConstantBuffers(1, Some(&blur_params)); + ctx.PSSetConstantBuffers(1, Some(&blur_params)); + ctx.PSSetSamplers(0, Some(slice::from_ref(&self.globals.sampler))); + ctx.PSSetShaderResources(0, Some(slice::from_ref(source_srv))); + ctx.OMSetBlendState(blend, None, 0xFFFFFFFF); + ctx.DrawInstanced(vertex_count, 1, 0, 0); + // Unbind the source so the target can be rebound as a render target next. + ctx.PSSetShaderResources(0, Some(&null_srv)); + } + Ok(()) + } + + /// Blur `source_srv` (full-resolution) using the half-res ping/pong textures and composite the + /// result into `target_rtv`, clipped to `bounds`/`corner_radii`/`content_mask` and modulated + /// by `opacity`. Shared by the backdrop and content-filter paths. + #[allow(clippy::too_many_arguments)] + fn dx_blur_and_composite( + &self, + source_srv: &Option, + target_rtv: &Option, + bounds: Bounds, + content_mask: Bounds, + corner_radii: [f32; 4], + blur_radius: f32, + opacity: f32, + // Backdrop clips to the rounded rect; content (`filter`) bleeds past its bounds. + clip_rounded: bool, + ) -> Result<()> { + // Sigma is halved because the blur runs at half resolution. + let sigma = (blur_radius * 0.5).max(0.0); + if sigma <= 0.0 { + return Ok(()); + } + // Span ±3σ. If that needs more than 32 taps, spread the taps apart (tap_step > 1) rather + // than truncating the kernel — keeps very large radii from clipping. Matches wgpu. + let ideal_taps = (3.0 * sigma).ceil(); + let tap_count = ideal_taps.clamp(1.0, 32.0); + let tap_step = (ideal_taps / tap_count).max(1.0); + // Content blur bleeds ~3·radius past the box, so its composite quad covers a dilated rect. + let composite_bounds = if clip_rounded { + bounds + } else { + bounds.dilate(ScaledPixels(3.0 * blur_radius)) + }; + let half_w = (self.width / 2).max(1); + let half_h = (self.height / 2).max(1); + let half_vp = D3D11_VIEWPORT { + TopLeftX: 0.0, + TopLeftY: 0.0, + Width: half_w as f32, + Height: half_h as f32, + MinDepth: 0.0, + MaxDepth: 1.0, + }; + let (full_vp, ping_rtv, ping_srv, pong_rtv, pong_srv) = { + let r = self.resources.as_ref().context("resources missing")?; + ( + r.viewport, + r.blur.ping_rtv.clone(), + r.blur.ping_srv.clone(), + r.blur.pong_rtv.clone(), + r.blur.pong_srv.clone(), + ) + }; + + // Downsample source -> ping, then separable gaussian ping -> pong -> ping. + self.dx_blur_pass( + &self.pipelines.blur_downsample_vertex, + &self.pipelines.blur_downsample_fragment, + &self.pipelines.blur_blend_replace, + &ping_rtv, + source_srv, + BlurParams { + downsample: 1.0, + ..Default::default() + }, + &half_vp, + D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST, + 3, + true, + )?; + self.dx_blur_pass( + &self.pipelines.blur_vertex, + &self.pipelines.blur_fragment, + &self.pipelines.blur_blend_replace, + &pong_rtv, + &ping_srv, + BlurParams { + direction: [1.0 / half_w as f32, 0.0], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + &half_vp, + D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST, + 3, + true, + )?; + self.dx_blur_pass( + &self.pipelines.blur_vertex, + &self.pipelines.blur_fragment, + &self.pipelines.blur_blend_replace, + &ping_rtv, + &pong_srv, + BlurParams { + direction: [0.0, 1.0 / half_h as f32], + sigma, + tap_count, + tap_step, + ..Default::default() + }, + &half_vp, + D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST, + 3, + true, + )?; + // Composite the blurred result into the target (preserving its contents). + self.dx_blur_pass( + &self.pipelines.blur_composite_vertex, + &self.pipelines.blur_composite_fragment, + &self.pipelines.blur_blend_composite, + target_rtv, + &ping_srv, + BlurParams { + bounds: composite_bounds, + content_mask, + corner_radii, + opacity, + clip_rounded: if clip_rounded { 1.0 } else { 0.0 }, + ..Default::default() + }, + &full_vp, + D3D_PRIMITIVE_TOPOLOGY_TRIANGLESTRIP, + 4, + false, + )?; + Ok(()) + } + + /// Copy the offscreen scene texture into the swapchain render target. + fn dx_blit( + &self, + source_srv: &Option, + target_rtv: &Option, + ) -> Result<()> { + let full_vp = self + .resources + .as_ref() + .context("resources missing")? + .viewport; + self.dx_blur_pass( + &self.pipelines.blur_downsample_vertex, + &self.pipelines.blur_downsample_fragment, + &self.pipelines.blur_blend_replace, + target_rtv, + source_srv, + BlurParams::default(), + &full_vp, + D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST, + 3, + true, + ) + } + pub(crate) fn gpu_specs(&self) -> Result { let devices = self.devices.as_ref().context("devices missing")?; let desc = unsafe { devices.adapter.GetDesc1() }?; @@ -783,6 +1213,7 @@ impl DirectXResources { viewport, ) = create_resources(devices, &swap_chain, width, height)?; set_rasterizer_state(&devices.device, &devices.device_context)?; + let blur = BlurResources::new(&devices.device, width, height)?; Ok(Self { swap_chain, @@ -792,6 +1223,7 @@ impl DirectXResources { path_intermediate_msaa_texture, path_intermediate_msaa_view, path_intermediate_srv, + blur, viewport, }) } @@ -818,6 +1250,7 @@ impl DirectXResources { self.path_intermediate_msaa_texture = path_intermediate_msaa_texture; self.path_intermediate_msaa_view = path_intermediate_msaa_view; self.path_intermediate_srv = path_intermediate_srv; + self.blur = BlurResources::new(&devices.device, width, height)?; self.viewport = viewport; Ok(()) } @@ -882,6 +1315,36 @@ impl DirectXRenderPipelines { create_blend_state(device)?, )?; + let blur_downsample_vertex = create_vertex_shader( + device, + RawShaderBytes::new(ShaderModule::BlurDownsample, ShaderTarget::Vertex)?.as_bytes(), + )?; + let blur_downsample_fragment = create_fragment_shader( + device, + RawShaderBytes::new(ShaderModule::BlurDownsample, ShaderTarget::Fragment)?.as_bytes(), + )?; + let blur_vertex = create_vertex_shader( + device, + RawShaderBytes::new(ShaderModule::Blur, ShaderTarget::Vertex)?.as_bytes(), + )?; + let blur_fragment = create_fragment_shader( + device, + RawShaderBytes::new(ShaderModule::Blur, ShaderTarget::Fragment)?.as_bytes(), + )?; + let blur_composite_vertex = create_vertex_shader( + device, + RawShaderBytes::new(ShaderModule::BlurComposite, ShaderTarget::Vertex)?.as_bytes(), + )?; + let blur_composite_fragment = create_fragment_shader( + device, + RawShaderBytes::new(ShaderModule::BlurComposite, ShaderTarget::Fragment)?.as_bytes(), + )?; + let blur_params_buffer = create_constant_buffer(device, std::mem::size_of::())?; + let blur_blend_replace = create_blend_state_no_blend(device)?; + // Premultiplied (One / InvSrcAlpha) — the composite outputs a premultiplied blurred sample; + // straight-alpha blending would darken the faded edges. + let blur_blend_composite = create_blend_state_for_path_sprite(device)?; + Ok(Self { shadow_pipeline, quad_pipeline, @@ -891,6 +1354,15 @@ impl DirectXRenderPipelines { mono_sprites, subpixel_sprites, poly_sprites, + blur_downsample_vertex, + blur_downsample_fragment, + blur_vertex, + blur_fragment, + blur_composite_vertex, + blur_composite_fragment, + blur_params_buffer, + blur_blend_replace, + blur_blend_composite, }) } } @@ -969,6 +1441,47 @@ struct GlobalParams { _pad: [u32; 3], } +/// Mirrors the `BlurParams` cbuffer (register b1) in `shaders.hlsl`. 80 bytes (a multiple of 16, +/// as constant buffers require). Updated per blur pass via `update_buffer`. +#[repr(C)] +#[derive(Clone, Copy)] +struct BlurParams { + bounds: Bounds, + content_mask: Bounds, + corner_radii: [f32; 4], + direction: [f32; 2], + sigma: f32, + opacity: f32, + tap_count: f32, + /// 1.0 clips the composite to the rounded rect (backdrop); 0.0 lets content blur bleed past + /// its bounds like CSS `filter: blur`. + clip_rounded: f32, + /// 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the + /// origin, so a stationary element blurs identically at every window size); 0.0 = 1:1 copy + /// (the scene blit, which must not downsample). Downsample pass only. + downsample: f32, + /// Spacing between taps in pixels (gaussian passes only); >1 lets `tap_count` taps span very + /// large radii without truncating the gaussian, matching the wgpu backend. + tap_step: f32, +} + +impl Default for BlurParams { + fn default() -> Self { + BlurParams { + bounds: Bounds::default(), + content_mask: Bounds::default(), + corner_radii: [0.0; 4], + direction: [0.0, 0.0], + sigma: 0.0, + opacity: 1.0, + tap_count: 0.0, + clip_rounded: 0.0, + downsample: 0.0, + tap_step: 0.0, + } + } +} + struct PipelineState { label: &'static str, vertex: ID3D11VertexShader, @@ -1267,6 +1780,16 @@ fn create_resources( } #[inline] +/// Flatten a `Corners` into the `[tl, tr, br, bl]` order expected by the blur composite shader. +fn corner_radii_array(corners: Corners) -> [f32; 4] { + [ + corners.top_left.0, + corners.top_right.0, + corners.bottom_right.0, + corners.bottom_left.0, + ] +} + fn create_render_target_and_its_view( swap_chain: &IDXGISwapChain1, device: &ID3D11Device, @@ -1310,6 +1833,45 @@ fn create_path_intermediate_texture( Ok((texture, Some(shader_resource_view.unwrap()))) } +/// Create a color texture usable as both a render target and a shader resource, returning both +/// views. Used for the blur offscreen targets. +#[inline] +fn create_color_target( + device: &ID3D11Device, + width: u32, + height: u32, +) -> Result<( + ID3D11Texture2D, + Option, + Option, +)> { + let texture = unsafe { + let mut output = None; + let desc = D3D11_TEXTURE2D_DESC { + Width: width.max(1), + Height: height.max(1), + MipLevels: 1, + ArraySize: 1, + Format: RENDER_TARGET_FORMAT, + SampleDesc: DXGI_SAMPLE_DESC { + Count: 1, + Quality: 0, + }, + Usage: D3D11_USAGE_DEFAULT, + BindFlags: (D3D11_BIND_RENDER_TARGET.0 | D3D11_BIND_SHADER_RESOURCE.0) as u32, + CPUAccessFlags: 0, + MiscFlags: 0, + }; + device.CreateTexture2D(&desc, None, Some(&mut output))?; + output.unwrap() + }; + let mut rtv = None; + unsafe { device.CreateRenderTargetView(&texture, None, Some(&mut rtv))? }; + let mut srv = None; + unsafe { device.CreateShaderResourceView(&texture, None, Some(&mut srv))? }; + Ok((texture, rtv, srv)) +} + #[inline] fn create_path_intermediate_msaa_texture_and_view( device: &ID3D11Device, @@ -1458,6 +2020,35 @@ fn create_blend_state_for_path_sprite(device: &ID3D11Device) -> Result Result { + let desc = D3D11_BUFFER_DESC { + ByteWidth: byte_size.next_multiple_of(16) as u32, + Usage: D3D11_USAGE_DYNAMIC, + BindFlags: D3D11_BIND_CONSTANT_BUFFER.0 as u32, + CPUAccessFlags: D3D11_CPU_ACCESS_WRITE.0 as u32, + ..Default::default() + }; + let mut buffer = None; + unsafe { device.CreateBuffer(&desc, None, Some(&mut buffer)) }?; + Ok(buffer.unwrap()) +} + +/// A blend state that overwrites the target (no blending) — used for the blur downsample and +/// gaussian passes. +#[inline] +fn create_blend_state_no_blend(device: &ID3D11Device) -> Result { + let mut desc = D3D11_BLEND_DESC::default(); + desc.RenderTarget[0].BlendEnable = false.into(); + desc.RenderTarget[0].RenderTargetWriteMask = D3D11_COLOR_WRITE_ENABLE_ALL.0 as u8; + unsafe { + let mut state = None; + device.CreateBlendState(&desc, Some(&mut state))?; + Ok(state.unwrap()) + } +} + #[inline] fn create_vertex_shader(device: &ID3D11Device, bytes: &[u8]) -> Result { unsafe { @@ -1604,6 +2195,9 @@ pub(crate) mod shader_resources { SubpixelSprite, PolychromeSprite, EmojiRasterization, + BlurDownsample, + Blur, + BlurComposite, } #[derive(Copy, Clone, Debug, Eq, PartialEq)] @@ -1681,6 +2275,18 @@ pub(crate) mod shader_resources { ShaderTarget::Vertex => EMOJI_RASTERIZATION_VERTEX_BYTES, ShaderTarget::Fragment => EMOJI_RASTERIZATION_FRAGMENT_BYTES, }, + ShaderModule::BlurDownsample => match target { + ShaderTarget::Vertex => BLUR_DOWNSAMPLE_VERTEX_BYTES, + ShaderTarget::Fragment => BLUR_DOWNSAMPLE_FRAGMENT_BYTES, + }, + ShaderModule::Blur => match target { + ShaderTarget::Vertex => BLUR_VERTEX_BYTES, + ShaderTarget::Fragment => BLUR_FRAGMENT_BYTES, + }, + ShaderModule::BlurComposite => match target { + ShaderTarget::Vertex => BLUR_COMPOSITE_VERTEX_BYTES, + ShaderTarget::Fragment => BLUR_COMPOSITE_FRAGMENT_BYTES, + }, }; Self { inner: bytes } } @@ -1768,6 +2374,9 @@ pub(crate) mod shader_resources { ShaderModule::SubpixelSprite => "subpixel_sprite", ShaderModule::PolychromeSprite => "polychrome_sprite", ShaderModule::EmojiRasterization => "emoji_rasterization", + ShaderModule::BlurDownsample => "blur_downsample", + ShaderModule::Blur => "blur", + ShaderModule::BlurComposite => "blur_composite", } } } diff --git a/crates/gpui_windows/src/shaders.hlsl b/crates/gpui_windows/src/shaders.hlsl index 89c12489b2..9e874d3a2f 100644 --- a/crates/gpui_windows/src/shaders.hlsl +++ b/crates/gpui_windows/src/shaders.hlsl @@ -1256,3 +1256,122 @@ float4 polychrome_sprite_fragment(PolychromeSpriteFragmentInput input): SV_Targe color.a *= sprite.opacity * saturate(0.5 - distance); return color; } + +/* +** +** Blur (backdrop-filter / filter) +** +** Shared by backdrop and content blur. The source texture is bound at t0 (t_sprite) and the +** parameters in the BlurParams constant buffer at b1. Three passes: downsample (full -> half +** res), separable gaussian (run twice), and a composite into a rounded rectangle. +*/ + +cbuffer BlurParams: register(b1) { + Bounds blur_bounds; + Bounds blur_content_mask; + float4 blur_corner_radii; + float2 blur_direction; + float blur_sigma; + float blur_opacity; + float blur_tap_count; + float blur_clip_rounded; + // 1.0 = snapped 2:1 box downsample (anchor the half-res grid to a fixed 2px grid at the origin + // so a stationary element blurs identically at every window size); 0.0 = 1:1 copy (scene blit). + float blur_downsample; + // Spacing between taps in pixels (gaussian passes only); >1 lets `blur_tap_count` taps span + // very large radii without truncating the gaussian. + float blur_tap_step; +}; + +struct BlurVertexOutput { + float4 position: SV_Position; + float2 uv: TEXCOORD0; +}; + +BlurVertexOutput blur_fullscreen(uint vertex_id) { + float2 uv = float2(float((vertex_id << 1u) & 2u), float(vertex_id & 2u)); + BlurVertexOutput output; + output.uv = uv; + output.position = float4(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, 0.0, 1.0); + return output; +} + +BlurVertexOutput blur_downsample_vertex(uint vertex_id: SV_VertexID) { + return blur_fullscreen(vertex_id); +} + +float4 blur_downsample_fragment(BlurVertexOutput input): SV_Target { + if (blur_downsample > 0.5) { + // Snapped 2:1 box downsample. Half-res texel `px` samples source at full-res coordinate + // 2*px + 1 (the boundary between source texels 2*px and 2*px+1), so one bilinear tap + // averages exactly that pair. Anchored to the origin and independent of the viewport size, + // so an element at fixed pixels blurs identically at every window size — otherwise the + // implicit floor(W/2) grid stretches and the halo wobbles by ~1px on resize. + uint sw, sh; + t_sprite.GetDimensions(sw, sh); + float2 src_uv = (floor(input.position.xy) * 2.0 + 1.0) / float2(sw, sh); + return t_sprite.SampleLevel(s_sprite, src_uv, 0.0); + } + // 1:1 copy at matching resolution (used to blit the offscreen scene into the swapchain). + return t_sprite.SampleLevel(s_sprite, input.uv, 0.0); +} + +BlurVertexOutput blur_vertex(uint vertex_id: SV_VertexID) { + return blur_fullscreen(vertex_id); +} + +float4 blur_fragment(BlurVertexOutput input): SV_Target { + int taps = int(blur_tap_count); + float4 color = float4(0.0, 0.0, 0.0, 0.0); + float weight_sum = 0.0; + [loop] + for (int i = -taps; i <= taps; i++) { + float offset = float(i) * blur_tap_step; + float weight = gaussian(offset, blur_sigma); + color += t_sprite.SampleLevel(s_sprite, input.uv + blur_direction * offset, 0.0) * weight; + weight_sum += weight; + } + return color / max(weight_sum, 1e-5); +} + +struct BlurCompositeVertexOutput { + float4 position: SV_Position; + float4 clip_distance: SV_ClipDistance; +}; + +struct BlurCompositeFragmentInput { + float4 position: SV_Position; +}; + +BlurCompositeVertexOutput blur_composite_vertex(uint vertex_id: SV_VertexID) { + float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + BlurCompositeVertexOutput output; + output.position = to_device_position(unit_vertex, blur_bounds); + output.clip_distance = distance_from_clip_rect(unit_vertex, blur_bounds, blur_content_mask); + return output; +} + +float4 blur_composite_fragment(BlurCompositeFragmentInput input): SV_Target { + // Sample the half-res blur by screen position, on the SAME fixed 2:1 grid the snapped downsample + // wrote (anchored at the origin, independent of viewport parity): 2 * the half-res texture size + // maps screen pixel p to half-res texel p/2 at every window size, so it doesn't wobble on resize. + uint hw, hh; + t_sprite.GetDimensions(hw, hh); + float2 uv = input.position.xy / (2.0 * float2(hw, hh)); + float4 blurred = t_sprite.SampleLevel(s_sprite, uv, 0.0); + Corners radii; + radii.top_left = blur_corner_radii.x; + radii.top_right = blur_corner_radii.y; + radii.bottom_right = blur_corner_radii.z; + radii.bottom_left = blur_corner_radii.w; + float distance = quad_sdf(input.position.xy, blur_bounds, radii); + // Backdrop clips to the rounded rect (the panel has a defined shape); content blur bleeds past + // its bounds like CSS `filter: blur`, so its shape comes from the blurred group's own alpha. + float coverage = blur_clip_rounded > 0.5 ? saturate(0.5 - distance) : 1.0; + // The blurred sample is premultiplied (blurring against the transparent surround scales rgb + // with the fading alpha), so output premultiplied and use a premultiplied-blend state. A + // backdrop's scene is opaque (so this replaces); a content-filter group is transparent outside + // its subtree (so the target shows through there instead of darkening). + float a = coverage * blur_opacity; + return float4(blurred.rgb * a, blurred.a * a); +}