Implement nested SVG clip paths in the D3D11 backend.

Partially addresses rust-windowing#372.

renderer/src/gpu_data.rs

@@ -179,8 +179,6 @@ pub struct SegmentIndicesD3D11 {
 #[derive(Clone, Debug)]
 pub struct ClippedPathInfo {
     /// The ID of the batch containing the clips.
-    /// 
-    /// In the current implementation, this is always 0.
     pub clip_batch_id: TileBatchId,
 
     /// The number of paths that have clips.
@@ -205,6 +203,12 @@ pub struct PathBatchIndex(pub u32);
 #[derive(Clone, Copy, PartialEq, Debug)]
 pub struct TileBatchId(pub u32);
 
+#[derive(Clone, Copy, PartialEq, Debug)]
+pub struct GlobalPathId {
+    pub batch_id: TileBatchId,
+    pub path_index: PathBatchIndex,
+}
+
 #[derive(Clone, Debug)]
 pub enum DrawTileBatch {
     D3D9(DrawTileBatchD3D9),
@@ -433,6 +437,13 @@ impl PathBatchIndex {
     }
 }
 
+impl GlobalPathId {
+    #[inline]
+    pub fn none() -> GlobalPathId {
+        GlobalPathId { batch_id: TileBatchId(!0), path_index: PathBatchIndex(!0) }
+    }
+}
+
 impl AlphaTileId {
     #[inline]
     pub fn new(next_alpha_tile_index: &[AtomicUsize; ALPHA_TILE_LEVEL_COUNT], level: usize) 

renderer/src/scene.rs

@@ -352,6 +352,7 @@ pub struct DrawPath {
 #[derive(Clone, Debug)]
 pub struct ClipPath {
     pub outline: Outline,
+    pub clip_path: Option<ClipPathId>,
     pub fill_rule: FillRule,
     pub name: String,
 }
@@ -447,14 +448,24 @@ impl DrawPath {
 impl ClipPath {
     #[inline]
     pub fn new(outline: Outline) -> ClipPath {
-        ClipPath { outline, fill_rule: FillRule::Winding, name: String::new() }
+        ClipPath { outline, clip_path: None, fill_rule: FillRule::Winding, name: String::new() }
     }
 
     #[inline]
     pub fn outline(&self) -> &Outline {
         &self.outline
     }
 
+    #[inline]
+    pub(crate) fn clip_path(&self) -> Option<ClipPathId> {
+        self.clip_path
+    }
+
+    #[inline]
+    pub fn set_clip_path(&mut self, new_clip_path: Option<ClipPathId>) {
+        self.clip_path = new_clip_path
+    }
+
     #[inline]
     pub(crate) fn fill_rule(&self) -> FillRule {
         self.fill_rule

renderer/src/tiler.rs

@@ -15,8 +15,8 @@ use crate::builder::{BuiltPath, BuiltPathBinCPUData, BuiltPathData, ObjectBuilde
 use crate::gpu::options::RendererLevel;
 use crate::gpu_data::AlphaTileId;
 use crate::options::PrepareMode;
-use crate::scene::PathId;
-use crate::tiles::{DrawTilingPathInfo, TILE_HEIGHT, TILE_WIDTH, TilingPathInfo};
+use crate::scene::{ClipPathId, PathId};
+use crate::tiles::{TILE_HEIGHT, TILE_WIDTH, TilingPathInfo};
 use pathfinder_content::fill::FillRule;
 use pathfinder_content::outline::{ContourIterFlags, Outline};
 use pathfinder_content::segment::Segment;
@@ -41,15 +41,14 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> {
                       fill_rule: FillRule,
                       view_box: RectF,
                       prepare_mode: &PrepareMode,
+                      clip_path_id: Option<ClipPathId>,
                       built_clip_paths: &'a [BuiltPath],
                       path_info: TilingPathInfo)
                       -> Tiler<'a, 'b, 'c, 'd> {
         let bounds = outline.bounds().intersection(view_box).unwrap_or(RectF::default());
 
-        let clip_path = match path_info {
-            TilingPathInfo::Draw(DrawTilingPathInfo { clip_path_id: Some(clip_path_id), .. }) => {
-                Some(&built_clip_paths[clip_path_id.0 as usize])
-            }
+        let clip_path = match clip_path_id {
+            Some(clip_path_id) => Some(&built_clip_paths[clip_path_id.0 as usize]),
             _ => None,
         };
 
@@ -58,6 +57,7 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> {
                                                 view_box,
                                                 fill_rule,
                                                 prepare_mode,
+                                                clip_path_id,
                                                 &path_info);
 
         Tiler { scene_builder, object_builder, outline, clip_path }
@@ -158,41 +158,6 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> {
 
             backdrops[column] += delta;
         }
-
-        /*
-
-        // Calculate clips.
-        let built_clip_path = match self.path_info {
-            TilingPathInfo::Draw(DrawTilingPathInfo {
-                built_clip_path: Some(built_clip_path),
-                ..
-            }) => built_clip_path,
-            _ => return,
-        };
-
-        let clip_tiles = self.object_builder
-                             .built_path
-                             .clip_tiles
-                             .as_mut()
-                             .expect("Where are the clip tiles?");
-
-        for draw_tile in &mut self.object_builder.built_path.tiles.data {
-            let tile_coords = vec2i(draw_tile.tile_x as i32, draw_tile.tile_y as i32);
-            let built_clip_tile = match built_clip_path.tiles.get(tile_coords) {
-                None => {
-                    draw_tile.alpha_tile_id = AlphaTileId(!0);
-                    continue;
-                }
-                Some(built_clip_tile) => built_clip_tile,
-            };
-
-            let clip_tile = clip_tiles.get_mut(tile_coords).unwrap();
-            clip_tile.dest_tile_id = draw_tile.alpha_tile_id;
-            clip_tile.dest_backdrop = draw_tile.backdrop as i32;
-            clip_tile.src_tile_id = built_clip_tile.alpha_tile_id;
-            clip_tile.src_backdrop = built_clip_tile.backdrop as i32;
-        }
-        */
     }
 }
 

renderer/src/tiles.rs

@@ -31,7 +31,6 @@ pub(crate) struct DrawTilingPathInfo {
     pub(crate) paint_id: PaintId,
     pub(crate) blend_mode: BlendMode,
     pub(crate) fill_rule: FillRule,
-    pub(crate) clip_path_id: Option<ClipPathId>,
 }
 
 impl TilingPathInfo {

resources/shaders/gl4/d3d11/propagate.cs.glsl

@@ -150,11 +150,6 @@ void main(){
                                                         clipTileRect,
                                                         clipTileCoord);
 
-
-
-
-
-
                 int thisClipAlphaTileIndex =
                     int(iClipTiles[clipTileIndex * 4 +
                                                                    2]<< 8)>> 8;
@@ -186,8 +181,6 @@ void main(){
 
                         drawTileBackdrop = 0;
                         needNewAlphaTile = false;
-                    } else {
-                        needNewAlphaTile = true;
                     }
                 }
             } else {

resources/shaders/metal/d3d11/propagate.cs.metal

@@ -61,7 +61,7 @@ uint calculateTileIndex(thread const uint& bufferOffset, thread const uint4& til
     return (bufferOffset + (tileCoord.y * (tileRect.z - tileRect.x))) + tileCoord.x;
 }
 
-kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bIndirectDrawParams& _303 [[buffer(6)]], device bAlphaTiles& _310 [[buffer(7)]], device bZBuffer& _381 [[buffer(10)]], device bFirstTileMap& _398 [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bIndirectDrawParams& _302 [[buffer(6)]], device bAlphaTiles& _309 [[buffer(7)]], device bZBuffer& _380 [[buffer(10)]], device bFirstTileMap& _397 [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
     uint columnIndex = gl_GlobalInvocationID.x;
     if (int(columnIndex) >= uColumnCount)
@@ -144,10 +144,6 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst
                         drawTileBackdrop = 0;
                         needNewAlphaTile = false;
                     }
-                    else
-                    {
-                        needNewAlphaTile = true;
-                    }
                 }
             }
             else
@@ -158,10 +154,10 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst
         }
         if (needNewAlphaTile)
         {
-            uint _306 = atomic_fetch_add_explicit((device atomic_uint*)&_303.iIndirectDrawParams[4], 1u, memory_order_relaxed);
-            uint drawBatchAlphaTileIndex = _306;
-            _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex;
-            _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 1u] = uint(clipAlphaTileIndex);
+            uint _305 = atomic_fetch_add_explicit((device atomic_uint*)&_302.iIndirectDrawParams[4], 1u, memory_order_relaxed);
+            uint drawBatchAlphaTileIndex = _305;
+            _309.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex;
+            _309.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 1u] = uint(clipAlphaTileIndex);
             drawAlphaTileIndex = int(drawBatchAlphaTileIndex) + uFirstAlphaTileIndex;
         }
         _175.iDrawTiles[(drawTileIndex * 4u) + 2u] = (uint(drawAlphaTileIndex) & 16777215u) | (uint(drawBackdropDelta) << uint(24));
@@ -170,12 +166,12 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst
         int tileMapIndex = (tileCoord_1.y * uFramebufferTileSize.x) + tileCoord_1.x;
         if ((zWrite && (drawTileBackdrop != 0)) && (drawAlphaTileIndex < 0))
         {
-            int _386 = atomic_fetch_max_explicit((device atomic_int*)&_381.iZBuffer[tileMapIndex], int(drawTileIndex), memory_order_relaxed);
+            int _385 = atomic_fetch_max_explicit((device atomic_int*)&_380.iZBuffer[tileMapIndex], int(drawTileIndex), memory_order_relaxed);
         }
         if ((drawTileBackdrop != 0) || (drawAlphaTileIndex >= 0))
         {
-            int _403 = atomic_exchange_explicit((device atomic_int*)&_398.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed);
-            int nextTileIndex = _403;
+            int _402 = atomic_exchange_explicit((device atomic_int*)&_397.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed);
+            int nextTileIndex = _402;
             _175.iDrawTiles[(drawTileIndex * 4u) + 0u] = uint(nextTileIndex);
         }
         currentBackdrop += drawBackdropDelta;

shaders/d3d11/propagate.cs.glsl

@@ -148,11 +148,6 @@ void main() {
                                                         clipTileRect,
                                                         clipTileCoord);
 
-/*
-                clipAlphaTileIndex =
-                    int(iClipTiles[clipTileIndex * 4 +
-                                   TILE_FIELD_BACKDROP_ALPHA_TILE_ID] << 8) >> 8;
-                                   */
                 int thisClipAlphaTileIndex =
                     int(iClipTiles[clipTileIndex * 4 +
                                    TILE_FIELD_BACKDROP_ALPHA_TILE_ID] << 8) >> 8;
@@ -166,8 +161,8 @@ void main() {
                         needNewAlphaTile = true;
                     } else {
                         if (drawTileBackdrop != 0) {
-                            // This is a solid draw tile, but there's a clip applied. Replace it with an
-                            // alpha tile pointing directly to the clip mask.
+                            // This is a solid draw tile, but there's a clip applied. Replace it
+                            // with an alpha tile pointing directly to the clip mask.
                             drawAlphaTileIndex = thisClipAlphaTileIndex;
                             clipAlphaTileIndex = -1;
                             needNewAlphaTile = false;
@@ -184,8 +179,6 @@ void main() {
                         // This is a blank clip tile. Cull the draw tile entirely.
                         drawTileBackdrop = 0;
                         needNewAlphaTile = false;
-                    } else {
-                        needNewAlphaTile = true;
                     }
                 }
             } else {

svg/src/lib.rs

@@ -40,7 +40,7 @@ const HAIRLINE_STROKE_WIDTH: f32 = 0.0333;
 pub struct SVGScene {
     pub scene: Scene,
     pub result_flags: BuildResultFlags,
-    pub clip_paths: HashMap<String, ClipPathId>,
+    pub clip_paths: HashMap<String, Outline>,
     gradients: HashMap<String, GradientInfo>,
 }
 
@@ -110,13 +110,12 @@ impl SVGScene {
                 }
 
                 if let Some(ref clip_path_name) = group.clip_path {
-                    if let Some(clip_path_id) = self.clip_paths.get(clip_path_name) {
-                        // TODO(pcwalton): Combine multiple clip paths if there's already one.
-                        if state.clip_path.is_some() {
-                            self.result_flags
-                                .insert(BuildResultFlags::UNSUPPORTED_MULTIPLE_CLIP_PATHS);
-                        }
-                        state.clip_path = Some(*clip_path_id);
+                    if let Some(clip_outline) = self.clip_paths.get(clip_path_name) {
+                        let mut clip_path = ClipPath::new((*clip_outline).clone());
+                        clip_path.set_clip_path(state.clip_path);
+                        clip_path.set_name(format!("ClipPath({})", clip_path_name));
+                        let clip_path_id = self.scene.push_clip_path(clip_path);
+                        state.clip_path = Some(clip_path_id);
                     }
                 }
 
@@ -187,13 +186,7 @@ impl SVGScene {
                     self.process_node(&kid, &state, &mut clip_outline);
                 }
 
-                if let Some(clip_outline) = clip_outline {
-                    // FIXME(pcwalton): Is the winding fill rule correct to use?
-                    let mut clip_path = ClipPath::new(clip_outline);
-                    clip_path.set_name(format!("ClipPath({})", node.id()));
-                    let clip_path_id = self.scene.push_clip_path(clip_path);
-                    self.clip_paths.insert(node.id().to_owned(), clip_path_id);
-                }
+                self.clip_paths.insert(node.id().to_owned(), clip_outline.unwrap());
             }
             NodeKind::Defs => {
                 // FIXME(pcwalton): This is wrong.