Filter API + Blur (#40)

Co-authored-by: Newspicel <newspicel+claude@pm.me>

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
 # ============================================================================

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::<Vec<_>>(),
+            )
+        }))
+}
+
+/// 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<Self>) -> 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");
+    });
+}

crates/gpui/src/bounds_tree.rs

@@ -28,6 +28,10 @@ where
     root: Option<usize>,
     /// Index of the leaf with the highest ordering (for fast-path lookups).
     max_leaf: Option<usize>,
+    /// 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<usize>,
     /// 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<U>) -> 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<U>) -> 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(),
         }

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<SubpixelSprite>,
     pub polychrome_sprites: Vec<PolychromeSprite>,
     pub surfaces: Vec<PaintSurface>,
+    pub backdrop_filters: Vec<BackdropFilter>,
+    pub filter_boundaries: Vec<FilterBoundary>,
 }
 
 #[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<Primitive>) {
         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<slice::Iter<'a, PolychromeSprite>>,
     surfaces_start: usize,
     surfaces_iter: Peekable<slice::Iter<'a, PaintSurface>>,
+    backdrop_filters_start: usize,
+    backdrop_filters_iter: Peekable<slice::Iter<'a, BackdropFilter>>,
+    filter_boundaries_start: usize,
+    filter_boundaries_iter: Peekable<slice::Iter<'a, FilterBoundary>>,
 }
 
 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<usize>,
     },
     Surfaces(Range<usize>),
+    BackdropFilters(Range<usize>),
+    /// 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<Shadow> 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<ScaledPixels>,
+    pub content_mask: ContentMask<ScaledPixels>,
+    pub corner_radii: Corners<ScaledPixels>,
+    /// 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<BackdropFilter> 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<ScaledPixels>,
+    pub content_mask: ContentMask<ScaledPixels>,
+    pub corner_radii: Corners<ScaledPixels>,
+    /// 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<FilterBoundary> 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<Pixels> {
         }
     }
 }
+
+#[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<ScaledPixels> {
+        Bounds {
+            origin: Point {
+                x: sp(0.0),
+                y: sp(0.0),
+            },
+            size: Size {
+                width: sp(100.0),
+                height: sp(100.0),
+            },
+        }
+    }
+
+    fn mask() -> ContentMask<ScaledPixels> {
+        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"]);
+    }
+}

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<BoxShadow>,
 
+    /// Filters applied to this element's own content and children (CSS `filter`).
+    pub filter: Vec<Filter>,
+
+    /// Filters applied to the content rendered behind this element (CSS `backdrop-filter`).
+    pub backdrop_filter: Vec<Filter>,
+
     /// 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(<px>)`.
+    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,

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(<radius>)`.
+    ///
+    /// 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.).<other_filter>()`). To replace the whole chain, use
+    /// [`Styled::filter`].
+    fn blur(mut self, radius: impl Into<Pixels>) -> 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<Vec<Filter>>) -> Self {
+        self.style().filter = Some(filters.into());
+        self
+    }
+
+    /// Blur the content rendered behind this element — a frosted-glass effect — like CSS
+    /// `backdrop-filter: blur(<radius>)`. 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<Pixels>) -> 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<Vec<Filter>>) -> 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 {

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<Pixels>,
+        corner_radii: Corners<Pixels>,
+        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<R>(
+        &mut self,
+        bounds: Bounds<Pixels>,
+        corner_radii: Corners<Pixels>,
+        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.

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<Mutex<InstanceBufferPool>>;
 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<Mutex<InstanceBufferPool>>,
@@ -132,9 +155,68 @@ pub(crate) struct MetalRenderer {
     core_video_texture_cache: core_video::metal_texture_cache::CVMetalTextureCache,
     path_intermediate_texture: Option<metal::Texture>,
     path_intermediate_msaa_texture: Option<metal::Texture>,
+    // 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<metal::Texture>,
+    blur_ping_texture: Option<metal::Texture>,
+    blur_pong_texture: Option<metal::Texture>,
+    /// 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<metal::Texture>,
     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<ScaledPixels>,
+    content_mask: Bounds<ScaledPixels>,
+    corner_radii: Corners<ScaledPixels>,
+    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<ScaledPixels>,
@@ -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<DevicePixels>,
+        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::<BlurUniform>() as u64,
+            &params 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<DevicePixels> as *const _,
+        );
+        encoder.set_fragment_bytes(
+            BlurInputIndex::Params as u64,
+            mem::size_of::<BlurUniform>() as u64,
+            &params 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<DevicePixels> 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<DevicePixels>,
+        bounds: Bounds<ScaledPixels>,
+        content_mask: Bounds<ScaledPixels>,
+        corner_radii: Corners<ScaledPixels>,
+        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<ScaledPixels>],
@@ -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;

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<float> source [[texture(0)]],
+    constant BlurParams &params [[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<float> source [[texture(0)]],
+    constant BlurParams &params [[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 &params [[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<float> source [[texture(0)]],
+    constant BlurParams &params [[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);
+}

crates/gpui_wgpu/src/shaders.wgsl

@@ -1348,3 +1348,136 @@ fn fs_surface(input: SurfaceVarying) -> @location(0) vec4<f32> {
 
     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<f32>,
+    direction: vec2<f32>,
+    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<uniform> blur_locals: BlurParams;
+@group(1) @binding(1) var t_blur: texture_2d<f32>;
+@group(1) @binding(2) var s_blur: sampler;
+
+struct BlurVarying {
+    @builtin(position) position: vec4<f32>,
+    @location(0) uv: vec2<f32>,
+    @location(3) clip_distances: vec4<f32>,
+}
+
+@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>(f32((vertex_id << 1u) & 2u), f32(vertex_id & 2u));
+    var out = BlurVarying();
+    out.uv = uv;
+    out.position = vec4<f32>(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, 0.0, 1.0);
+    out.clip_distances = vec4<f32>(1.0);
+    return out;
+}
+
+@fragment
+fn fs_blur_downsample(input: BlurVarying) -> @location(0) vec4<f32> {
+    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<f32> {
+    let sigma = blur_locals.sigma;
+    let taps = i32(blur_locals.tap_count);
+    let step = blur_locals.tap_step;
+    var color = vec4<f32>(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>(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<f32> {
+    if (any(input.clip_distances < vec4<f32>(0.0))) {
+        return vec4<f32>(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<f32>(blurred.rgb * c, blurred.a * c);
+}

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<wgpu::RenderPipeline>,
     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<wgpu::TextureView>,
     path_msaa_texture: Option<wgpu::Texture>,
     path_msaa_view: Option<wgpu::TextureView>,
+    /// 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<wgpu::Texture>,
+    scene_color_view: Option<wgpu::TextureView>,
+    /// Half-resolution ping/pong targets for the downsample + separable gaussian passes.
+    blur_ping_texture: Option<wgpu::Texture>,
+    blur_ping_view: Option<wgpu::TextureView>,
+    blur_pong_texture: Option<wgpu::Texture>,
+    blur_pong_view: Option<wgpu::TextureView>,
+    /// 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<wgpu::Texture>,
+    group_views: Vec<wgpu::TextureView>,
 }
 
 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<u64>,
     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::<BlurParams>() 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::<GlobalParams>() as u64;
         let gamma_size = std::mem::size_of::<GammaParams>() 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::<BlurParams>() 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: &current_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: &current_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, &current_target);
+                            }
+                            pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                                label: Some("main_pass_continued"),
+                                color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                                    view: &current_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: &current_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,
+                                        &current_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: &current_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(&params),
+        );
+        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::<BlurParams>() 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<ScaledPixels>,
+        content_mask: Bounds<ScaledPixels>,
+        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],

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);

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<ID3D11RenderTargetView>,
 }
 
 /// Direct3D objects
@@ -77,10 +101,68 @@ struct DirectXResources {
     path_intermediate_msaa_texture: ID3D11Texture2D,
     path_intermediate_msaa_view: Option<ID3D11RenderTargetView>,
 
+    // 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<ID3D11RenderTargetView>,
+    scene_color_srv: Option<ID3D11ShaderResourceView>,
+    ping: ID3D11Texture2D,
+    ping_rtv: Option<ID3D11RenderTargetView>,
+    ping_srv: Option<ID3D11ShaderResourceView>,
+    pong: ID3D11Texture2D,
+    pong_rtv: Option<ID3D11RenderTargetView>,
+    pong_srv: Option<ID3D11ShaderResourceView>,
+    // Kept alive for the lifetime of their views; indexed by isolation depth.
+    groups: Vec<ID3D11Texture2D>,
+    group_rtvs: Vec<Option<ID3D11RenderTargetView>>,
+    group_srvs: Vec<Option<ID3D11ShaderResourceView>>,
+}
+
+impl BlurResources {
+    fn new(device: &ID3D11Device, width: u32, height: u32) -> Result<Self> {
+        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<Shadow>,
     quad_pipeline: PipelineState<Quad>,
@@ -90,6 +172,18 @@ struct DirectXRenderPipelines {
     mono_sprites: PipelineState<MonochromeSprite>,
     subpixel_sprites: PipelineState<SubpixelSprite>,
     poly_sprites: PipelineState<PolychromeSprite>,
+    // 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<ID3D11RenderTargetView>,
+            Option<ID3D11ShaderResourceView>,
+            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(
+                                &current_srv,
+                                &current_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(&current_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(&current_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(
+                                &current_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(&current_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<ID3D11RenderTargetView>,
+        source_srv: &Option<ID3D11ShaderResourceView>,
+        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<ID3D11ShaderResourceView>; 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<ID3D11ShaderResourceView>,
+        target_rtv: &Option<ID3D11RenderTargetView>,
+        bounds: Bounds<ScaledPixels>,
+        content_mask: Bounds<ScaledPixels>,
+        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<ID3D11ShaderResourceView>,
+        target_rtv: &Option<ID3D11RenderTargetView>,
+    ) -> 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<GpuSpecs> {
         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::<BlurParams>())?;
+        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<ScaledPixels>,
+    content_mask: Bounds<ScaledPixels>,
+    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<T> {
     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<ScaledPixels>) -> [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<ID3D11RenderTargetView>,
+    Option<ID3D11ShaderResourceView>,
+)> {
+    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<ID3D11Ble
     }
 }
 
+/// Create a CPU-writable dynamic constant buffer of the given byte size (rounded up to 16).
+#[inline]
+fn create_constant_buffer(device: &ID3D11Device, byte_size: usize) -> Result<ID3D11Buffer> {
+    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<ID3D11BlendState> {
+    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<ID3D11VertexShader> {
     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",
             }
         }
     }

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