Auto merge of #380 - pcwalton:nested-clips, r=pcwalton

Implement nested SVG clip paths in the D3D11 backend.

Partially addresses #372.

renderer/src/builder.rs

@@ -14,11 +14,10 @@ use crate::concurrent::executor::Executor;
 use crate::gpu::blend::BlendModeExt;
 use crate::gpu::options::RendererLevel;
 use crate::gpu_data::{AlphaTileId, BackdropInfoD3D11, Clip, ClippedPathInfo, DiceMetadataD3D11};
-use crate::gpu_data::{DrawTileBatch, DrawTileBatchD3D9, DrawTileBatchD3D11, Fill, PathBatchIndex};
+use crate::gpu_data::{DrawTileBatch, DrawTileBatchD3D9, DrawTileBatchD3D11, Fill, GlobalPathId};
-use crate::gpu_data::{PathSource, PrepareTilesInfoD3D11};
+use crate::gpu_data::{PathBatchIndex, PathSource, PrepareTilesInfoD3D11, PropagateMetadataD3D11};
-use crate::gpu_data::{PropagateMetadataD3D11, RenderCommand, SegmentIndicesD3D11, SegmentsD3D11};
+use crate::gpu_data::{RenderCommand, SegmentIndicesD3D11, SegmentsD3D11, TileBatchDataD3D11};
-use crate::gpu_data::{TileBatchDataD3D11, TileBatchId, TileBatchTexture};
+use crate::gpu_data::{TileBatchId, TileBatchTexture, TileObjectPrimitive, TilePathInfoD3D11};
-use crate::gpu_data::{TileObjectPrimitive, TilePathInfoD3D11};
 use crate::options::{PrepareMode, PreparedBuildOptions, PreparedRenderTransform};
 use crate::paint::{PaintId, PaintInfo, PaintMetadata};
 use crate::scene::{ClipPathId, DisplayItem, DrawPath, DrawPathId, LastSceneInfo, PathId};
@@ -48,6 +47,8 @@ pub(crate) const ALPHA_TILES_PER_LEVEL: usize = 1 << (32 - ALPHA_TILE_LEVEL_COUN
 const CURVE_IS_QUADRATIC: u32 = 0x80000000;
 const CURVE_IS_CUBIC:     u32 = 0x40000000;
 
+const MAX_CLIP_BATCHES: u32 = 32;
+
 pub(crate) struct SceneBuilder<'a, 'b, 'c, 'd> {
     scene: &'a mut Scene,
     built_options: &'b PreparedBuildOptions,
@@ -265,6 +266,7 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> {
                                    path_object.fill_rule(),
                                    view_box,
                                    &prepare_mode,
+                                   path_object.clip_path(),
                                    &[],
                                    TilingPathInfo::Clip);
 
@@ -298,11 +300,11 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> {
                                    path_object.fill_rule(),
                                    view_box,
                                    &prepare_mode,
+                                   path_object.clip_path(),
                                    &built_clip_paths,
                                    TilingPathInfo::Draw(DrawTilingPathInfo {
             paint_id,
             blend_mode: path_object.blend_mode(),
-            clip_path_id: path_object.clip_path(),
             fill_rule: path_object.fill_rule(),
         }));
 
@@ -322,7 +324,7 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> {
                           paint_metadata: &[PaintMetadata],
                           prepare_mode: &PrepareMode,
                           built_paths: Option<BuiltPaths>) {
-        let mut tile_batch_builder = TileBatchBuilder::new(&prepare_mode, built_paths);
+        let mut tile_batch_builder = TileBatchBuilder::new(built_paths);
 
         // Prepare display items.
         for display_item in self.scene.display_list() {
@@ -410,6 +412,7 @@ impl BuiltPath {
            view_box_bounds: RectF,
            fill_rule: FillRule,
            prepare_mode: &PrepareMode,
+           clip_path_id: Option<ClipPathId>,
            tiling_path_info: &TilingPathInfo)
            -> BuiltPath {
         let paint_id = match *tiling_path_info {
@@ -427,13 +430,6 @@ impl BuiltPath {
 
         let tile_bounds = tiles::round_rect_out_to_tile_bounds(tile_map_bounds);
 
-        let clip_path_id = match *tiling_path_info {
-            TilingPathInfo::Draw(ref draw_tiling_path_info) => {
-                draw_tiling_path_info.clip_path_id
-            }
-            _ => None,
-        };
-
         let data = match *prepare_mode {
             PrepareMode::CPU => {
                 BuiltPathData::CPU(BuiltPathBinCPUData {
@@ -492,6 +488,7 @@ impl ObjectBuilder {
                       view_box_bounds: RectF,
                       fill_rule: FillRule,
                       prepare_mode: &PrepareMode,
+                      clip_path_id: Option<ClipPathId>,
                       tiling_path_info: &TilingPathInfo)
                       -> ObjectBuilder {
         let built_path = BuiltPath::new(path_id,
@@ -499,6 +496,7 @@ impl ObjectBuilder {
                                         view_box_bounds,
                                         fill_rule,
                                         prepare_mode,
+                                        clip_path_id,
                                         tiling_path_info);
         ObjectBuilder { built_path, bounds: path_bounds, fills: vec![] }
     }
@@ -650,7 +648,7 @@ impl TileBatchDataD3D11 {
     fn push(&mut self,
             path: &BuiltPath,
             global_path_id: PathId,
-            batch_clip_path_index: Option<PathBatchIndex>,
+            batch_clip_path_id: Option<GlobalPathId>,
             z_write: bool,
             sink: &SceneSink)
             -> PathBatchIndex {
@@ -662,7 +660,10 @@ impl TileBatchDataD3D11 {
             tile_offset: self.tile_count,
             path_index: batch_path_index,
             z_write: z_write as u32,
-            clip_path_index: batch_clip_path_index.unwrap_or(PathBatchIndex::none()),
+            clip_path_index: match batch_clip_path_id {
+                None => PathBatchIndex::none(),
+                Some(batch_clip_path_id) => batch_clip_path_id.path_index,
+            },
             backdrop_offset: self.prepare_info.backdrops.len() as u32,
             pad0: 0,
             pad1: 0,
@@ -716,13 +717,14 @@ impl TileBatchDataD3D11 {
 
         // Handle clip.
 
-        if batch_clip_path_index.is_none() {
+        let clip_batch_id = match batch_clip_path_id {
-            return batch_path_index;
+            None => return batch_path_index,
-        }
+            Some(batch_clip_path_id) => batch_clip_path_id.batch_id,
+        };
 
         if self.clipped_path_info.is_none() {
             self.clipped_path_info = Some(ClippedPathInfo {
-                clip_batch_id: TileBatchId(0),
+                clip_batch_id,
                 clipped_path_count: 0,
                 max_clipped_tile_count: 0,
                 clips: match sink.renderer_level {
@@ -836,31 +838,33 @@ impl SegmentsD3D11 {
 struct TileBatchBuilder {
     prepare_commands: Vec<RenderCommand>,
     draw_commands: Vec<RenderCommand>,
-    clip_id_to_path_batch_index: FxHashMap<ClipPathId, PathBatchIndex>,
+    clip_batches_d3d11: Option<ClipBatchesD3D11>,
     next_batch_id: TileBatchId,
     level: TileBatchBuilderLevel,
 }
 
 enum TileBatchBuilderLevel {
     D3D9 { built_paths: BuiltPaths },
-    D3D11 { clip_prepare_batch: TileBatchDataD3D11 },
+    D3D11,
 }
 
 impl TileBatchBuilder {
-    fn new(prepare_mode: &PrepareMode, built_paths: Option<BuiltPaths>) -> TileBatchBuilder {
+    fn new(built_paths: Option<BuiltPaths>) -> TileBatchBuilder {
         TileBatchBuilder {
             prepare_commands: vec![],
             draw_commands: vec![],
-            next_batch_id: TileBatchId(1),
+            next_batch_id: TileBatchId(MAX_CLIP_BATCHES),
-            clip_id_to_path_batch_index: FxHashMap::default(),
+            clip_batches_d3d11: match built_paths {
-            level: match built_paths {
                 None => {
-                    TileBatchBuilderLevel::D3D11 {
+                    Some(ClipBatchesD3D11 {
-                        clip_prepare_batch: TileBatchDataD3D11::new(TileBatchId(0),
+                        prepare_batches: vec![],
-                                                                    &prepare_mode,
+                        clip_id_to_path_batch_index: FxHashMap::default(),
-                                                                    PathSource::Clip),
+                    })
-                    }
                 }
+                Some(_) => None,
+            },
+            level: match built_paths {
+                None => TileBatchBuilderLevel::D3D11,
                 Some(built_paths) => TileBatchBuilderLevel::D3D9 { built_paths },
             },
         }
@@ -945,34 +949,16 @@ impl TileBatchBuilder {
             }
 
             // Add clip path if necessary.
-            let clip_path = match draw_path.clip_path_id {
+            let clip_path = match self.clip_batches_d3d11 {
                 None => None,
-                Some(clip_path_id) => {
+                Some(ref mut clip_batches_d3d11) => {
-                    match self.clip_id_to_path_batch_index.get(&clip_path_id) {
+                    add_clip_path_to_batch(scene,
-                        Some(&clip_path_batch_index) => Some(clip_path_batch_index),
+                                           sink,
-                        None => {
+                                           built_options,
-                            match self.level {
+                                           draw_path.clip_path_id,
-                                TileBatchBuilderLevel::D3D9 { .. } => None,
+                                           prepare_mode,
-                                TileBatchBuilderLevel::D3D11 { ref mut clip_prepare_batch } => {
+                                           0,
-                                    let clip_path =
+                                           clip_batches_d3d11)
-                                        prepare_clip_path_for_gpu_binning(scene,
-                                                                          built_options,
-                                                                          clip_path_id,
-                                                                          prepare_mode);
-                                    let clip_path_index =
-                                        clip_prepare_batch.push(&clip_path,
-                                                                clip_path_id.to_path_id(),
-                                                                None,
-                                                                true,
-                                                                sink);
-                                    self.clip_id_to_path_batch_index.insert(clip_path_id,
-                                                                            clip_path_index);
-                                    Some(clip_path_index)
-                                }
-                            }
-
-                        }
-                    }
                 }
             };
 
@@ -1071,19 +1057,21 @@ impl TileBatchBuilder {
                                         effective_view_box,
                                         draw_path.fill_rule(),
                                         &prepare_mode,
+                                        draw_path.clip_path(),
                                         &TilingPathInfo::Draw(DrawTilingPathInfo {
                                             paint_id,
                                             blend_mode: draw_path.blend_mode(),
-                                            clip_path_id: draw_path.clip_path(),
                                             fill_rule: draw_path.fill_rule(),
                                         }));
         Some(BuiltDrawPath::new(built_path, draw_path, paint_metadata))
     }
 
     fn send_to(self, sink: &SceneSink) {
-        if let TileBatchBuilderLevel::D3D11 { clip_prepare_batch } = self.level {
+        if let Some(clip_batches_d3d11) = self.clip_batches_d3d11 {
-            if clip_prepare_batch.path_count > 0 {
+            for prepare_batch in clip_batches_d3d11.prepare_batches.into_iter().rev() {
-                sink.listener.send(RenderCommand::PrepareClipTilesD3D11(clip_prepare_batch));
+                if prepare_batch.path_count > 0 {
+                    sink.listener.send(RenderCommand::PrepareClipTilesD3D11(prepare_batch));
+                }
             }
         }
 
@@ -1096,11 +1084,76 @@ impl TileBatchBuilder {
     }
 }
 
+struct ClipBatchesD3D11 {
+    // Will be submitted in reverse (LIFO) order.
+    prepare_batches: Vec<TileBatchDataD3D11>,
+    clip_id_to_path_batch_index: FxHashMap<ClipPathId, PathBatchIndex>,
+}
+
+fn add_clip_path_to_batch(scene: &Scene,
+                          sink: &SceneSink,
+                          built_options: &PreparedBuildOptions,
+                          clip_path_id: Option<ClipPathId>,
+                          prepare_mode: &PrepareMode,
+                          clip_level: usize,
+                          clip_batches_d3d11: &mut ClipBatchesD3D11)
+                          -> Option<GlobalPathId> {
+    match clip_path_id {
+        None => None,
+        Some(clip_path_id) => {
+            match clip_batches_d3d11.clip_id_to_path_batch_index.get(&clip_path_id) {
+                Some(&clip_path_batch_index) => {
+                    Some(GlobalPathId {
+                        batch_id: TileBatchId(clip_level as u32),
+                        path_index: clip_path_batch_index,
+                    })
+                }
+                None => {
+                    let PreparedClipPath {
+                        built_path: clip_path,
+                        subclip_id
+                    } = prepare_clip_path_for_gpu_binning(scene,
+                                                          sink,
+                                                          built_options,
+                                                          clip_path_id,
+                                                          prepare_mode,
+                                                          clip_level,
+                                                          clip_batches_d3d11);
+                    while clip_level >= clip_batches_d3d11.prepare_batches.len() {
+                        let clip_tile_batch_id =
+                            TileBatchId(clip_batches_d3d11.prepare_batches.len() as u32);
+                        clip_batches_d3d11.prepare_batches
+                                          .push(TileBatchDataD3D11::new(clip_tile_batch_id,
+                                                                        &prepare_mode,
+                                                                        PathSource::Clip));
+                    }
+                    let clip_path_batch_index =
+                        clip_batches_d3d11.prepare_batches[clip_level]
+                                          .push(&clip_path,
+                                                clip_path_id.to_path_id(),
+                                                subclip_id,
+                                                true,
+                                                sink);
+                    clip_batches_d3d11.clip_id_to_path_batch_index.insert(clip_path_id,
+                                                                          clip_path_batch_index);
+                    Some(GlobalPathId {
+                        batch_id: TileBatchId(clip_level as u32),
+                        path_index: clip_path_batch_index,
+                    })
+                }
+            }
+        }
+    }
+}
+
 fn prepare_clip_path_for_gpu_binning(scene: &Scene,
+                                     sink: &SceneSink,
                                      built_options: &PreparedBuildOptions,
                                      clip_path_id: ClipPathId,
-                                     prepare_mode: &PrepareMode)
+                                     prepare_mode: &PrepareMode,
-                                     -> BuiltPath {
+                                     clip_level: usize,
+                                     clip_batches_d3d11: &mut ClipBatchesD3D11)
+                                     -> PreparedClipPath {
     let transform = match *prepare_mode {
         PrepareMode::GPU { transform } => transform,
         PrepareMode::CPU | PrepareMode::TransformCPUBinGPU => {
@@ -1109,14 +1162,34 @@ fn prepare_clip_path_for_gpu_binning(scene: &Scene,
     };
     let effective_view_box = scene.effective_view_box(built_options);
     let clip_path = scene.get_clip_path(clip_path_id);
+
+    // Add subclip path if necessary.
+    let subclip_id = add_clip_path_to_batch(scene,
+                                            sink,
+                                            built_options,
+                                            clip_path.clip_path(),
+                                            prepare_mode,
+                                            clip_level + 1,
+                                            clip_batches_d3d11);
+
     let path_bounds = transform * clip_path.outline().bounds();
+
     // TODO(pcwalton): Clip to view box!
-    BuiltPath::new(clip_path_id.to_path_id(),
+
-                   path_bounds,
+    let built_path = BuiltPath::new(clip_path_id.to_path_id(),
-                   effective_view_box,
+                                    path_bounds,
-                   clip_path.fill_rule(),
+                                    effective_view_box,
-                   &prepare_mode,
+                                    clip_path.fill_rule(),
-                   &TilingPathInfo::Clip)
+                                    &prepare_mode,
+                                    clip_path.clip_path(),
+                                    &TilingPathInfo::Clip);
+
+    PreparedClipPath { built_path, subclip_id }
+}
+
+struct PreparedClipPath {
+    built_path: BuiltPath,
+    subclip_id: Option<GlobalPathId>,
 }
 
 fn fixup_batch_for_new_path_if_possible(batch_color_texture: &mut Option<TileBatchTexture>,

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,7 +448,7 @@ 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]
@@ -455,6 +456,16 @@ impl ClipPath {
         &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::scene::{ClipPathId, PathId};
-use crate::tiles::{DrawTilingPathInfo, TILE_HEIGHT, TILE_WIDTH, TilingPathInfo};
+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 {
+        let clip_path = match clip_path_id {
-            TilingPathInfo::Draw(DrawTilingPathInfo { clip_path_id: Some(clip_path_id), .. }) => {
+            Some(clip_path_id) => Some(&built_clip_paths[clip_path_id.0 as usize]),
-                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 _305 = atomic_fetch_add_explicit((device atomic_uint*)&_302.iIndirectDrawParams[4], 1u, memory_order_relaxed);
-            uint drawBatchAlphaTileIndex = _306;
+            uint drawBatchAlphaTileIndex = _305;
-            _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex;
+            _309.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex;
-            _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 1u] = uint(clipAlphaTileIndex);
+            _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 _402 = atomic_exchange_explicit((device atomic_int*)&_397.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed);
-            int nextTileIndex = _403;
+            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
+                            // This is a solid draw tile, but there's a clip applied. Replace it
-                            // alpha tile pointing directly to the clip mask.
+                            // 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) {
+                    if let Some(clip_outline) = self.clip_paths.get(clip_path_name) {
-                        // TODO(pcwalton): Combine multiple clip paths if there's already one.
+                        let mut clip_path = ClipPath::new((*clip_outline).clone());
-                        if state.clip_path.is_some() {
+                        clip_path.set_clip_path(state.clip_path);
-                            self.result_flags
+                        clip_path.set_name(format!("ClipPath({})", clip_path_name));
-                                .insert(BuildResultFlags::UNSUPPORTED_MULTIPLE_CLIP_PATHS);
+                        let clip_path_id = self.scene.push_clip_path(clip_path);
-                        }
+                        state.clip_path = Some(clip_path_id);
-                        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 {
+                self.clip_paths.insert(node.id().to_owned(), clip_outline.unwrap());
-                    // 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);
-                }
             }
             NodeKind::Defs => {
                 // FIXME(pcwalton): This is wrong.