Implement the infrastructure needed to support multiple clip paths.

This also lays the groundwork needed to reduce batch breaks between solid and
alpha tiles.

Cargo.lock

@@ -838,6 +838,14 @@ name = "fuchsia-zircon-sys"
 version = "0.3.3"
 source = "registry+https://github.com/rust-lang/crates.io-index"
 
+[[package]]
+name = "fxhash"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+dependencies = [
+ "byteorder 1.3.4 (registry+https://github.com/rust-lang/crates.io-index)",
+]
+
 [[package]]
 name = "gcc"
 version = "0.3.55"
@@ -1725,6 +1733,7 @@ dependencies = [
  "bitflags 1.2.1 (registry+https://github.com/rust-lang/crates.io-index)",
  "byteorder 1.3.4 (registry+https://github.com/rust-lang/crates.io-index)",
  "crossbeam-channel 0.4.2 (registry+https://github.com/rust-lang/crates.io-index)",
+ "fxhash 0.2.1 (registry+https://github.com/rust-lang/crates.io-index)",
  "half 1.5.0 (registry+https://github.com/rust-lang/crates.io-index)",
  "hashbrown 0.7.1 (registry+https://github.com/rust-lang/crates.io-index)",
  "instant 0.1.2 (registry+https://github.com/rust-lang/crates.io-index)",
@@ -2839,6 +2848,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
 "checksum fs_extra 1.1.0 (registry+https://github.com/rust-lang/crates.io-index)" = "5f2a4a2034423744d2cc7ca2068453168dcdb82c438419e639a26bd87839c674"
 "checksum fuchsia-zircon 0.3.3 (registry+https://github.com/rust-lang/crates.io-index)" = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82"
 "checksum fuchsia-zircon-sys 0.3.3 (registry+https://github.com/rust-lang/crates.io-index)" = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7"
+"checksum fxhash 0.2.1 (registry+https://github.com/rust-lang/crates.io-index)" = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c"
 "checksum gcc 0.3.55 (registry+https://github.com/rust-lang/crates.io-index)" = "8f5f3913fa0bfe7ee1fd8248b6b9f42a5af4b9d65ec2dd2c3c26132b950ecfc2"
 "checksum getrandom 0.1.14 (registry+https://github.com/rust-lang/crates.io-index)" = "7abc8dd8451921606d809ba32e95b6111925cd2906060d2dcc29c070220503eb"
 "checksum gif 0.10.3 (registry+https://github.com/rust-lang/crates.io-index)" = "471d90201b3b223f3451cd4ad53e34295f16a1df17b1edf3736d47761c3981af"

renderer/Cargo.toml

@@ -12,6 +12,7 @@ homepage = "https://github.com/servo/pathfinder"
 bitflags = "1.0"
 byteorder = "1.2"
 crossbeam-channel = "0.4"
+fxhash = "0.2"
 half = "1.5"
 hashbrown = "0.7"
 rayon = "1.0"

renderer/src/builder.rs

@@ -12,8 +12,8 @@
 
 use crate::concurrent::executor::Executor;
 use crate::gpu::renderer::{BlendModeExt, MASK_TILES_ACROSS, MASK_TILES_DOWN};
-use crate::gpu_data::{AlphaTileId, Fill, FillBatchEntry, RenderCommand, Tile, TileBatch};
+use crate::gpu_data::{AlphaTileId, Clip, ClipBatch, ClipBatchKey, ClipBatchKind, Fill, FillBatchEntry, RenderCommand};
-use crate::gpu_data::{TileBatchTexture, TileObjectPrimitive};
+use crate::gpu_data::{Tile, TileBatch, TileBatchTexture, TileObjectPrimitive};
 use crate::options::{PreparedBuildOptions, PreparedRenderTransform, RenderCommandListener};
 use crate::paint::{PaintInfo, PaintMetadata};
 use crate::scene::{DisplayItem, Scene};
@@ -31,14 +31,17 @@ use pathfinder_geometry::util;
 use pathfinder_geometry::vector::{Vector2F, Vector2I, vec2f, vec2i};
 use pathfinder_gpu::TextureSamplingFlags;
 use pathfinder_simd::default::{F32x4, I32x4};
-use std::sync::atomic::{AtomicUsize, Ordering};
+use std::sync::atomic::AtomicUsize;
 use instant::Instant;
 use std::u32;
 
+pub(crate) const ALPHA_TILE_LEVEL_COUNT: usize = 2;
+pub(crate) const ALPHA_TILES_PER_LEVEL: usize = 1 << (32 - ALPHA_TILE_LEVEL_COUNT + 1);
+
 pub(crate) struct SceneBuilder<'a, 'b> {
     scene: &'a mut Scene,
     built_options: &'b PreparedBuildOptions,
-    next_alpha_tile_index: AtomicUsize,
+    next_alpha_tile_indices: [AtomicUsize; ALPHA_TILE_LEVEL_COUNT],
     pub(crate) listener: Box<dyn RenderCommandListener>,
 }
 
@@ -57,7 +60,6 @@ struct BuiltDrawPath {
     color_texture: Option<TileBatchTexture>,
     sampling_flags_1: TextureSamplingFlags,
     mask_0_fill_rule: FillRule,
-    mask_1_fill_rule: Option<FillRule>,
 }
 
 #[derive(Debug)]
@@ -65,7 +67,7 @@ pub(crate) struct BuiltPath {
     pub solid_tiles: SolidTiles,
     pub empty_tiles: Vec<BuiltTile>,
     pub single_mask_tiles: Vec<BuiltTile>,
-    pub dual_mask_tiles: Vec<BuiltTile>,
+    pub clip_tiles: Vec<BuiltClip>,
     pub tiles: DenseTileMap<TileObjectPrimitive>,
     pub fill_rule: FillRule,
 }
@@ -76,6 +78,12 @@ pub struct BuiltTile {
     pub tile: Tile,
 }
 
+#[derive(Clone, Copy, Debug)]
+pub struct BuiltClip {
+    pub clip: Clip,
+    pub key: ClipBatchKey,
+}
+
 #[derive(Clone, Debug)]
 pub(crate) enum SolidTiles {
     Occluders(Vec<Occluder>),
@@ -96,7 +104,7 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
         SceneBuilder {
             scene,
             built_options,
-            next_alpha_tile_index: AtomicUsize::new(0),
+            next_alpha_tile_indices: [AtomicUsize::new(0), AtomicUsize::new(0)],
             listener,
         }
     }
@@ -216,7 +224,6 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
             color_texture: paint_metadata.tile_batch_texture(),
             sampling_flags_1: TextureSamplingFlags::empty(),
             mask_0_fill_rule: path_object.fill_rule(),
-            mask_1_fill_rule: built_clip_path.map(|_| FillRule::Winding),
         }
     }
 
@@ -226,6 +233,29 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
         }
     }
 
+    fn build_clips(&self, built_draw_paths: &[BuiltDrawPath]) {
+        let mut built_clip_tiles = vec![];
+        for built_draw_path in built_draw_paths {
+            for built_clip_tile in &built_draw_path.path.clip_tiles {
+                built_clip_tiles.push(*built_clip_tile);
+            }
+        }
+
+        built_clip_tiles.sort_by_key(|built_clip_tile| built_clip_tile.key);
+
+        let mut batches: Vec<ClipBatch> = vec![];
+        for built_clip_tile in built_clip_tiles {
+            if batches.is_empty() || batches.last_mut().unwrap().key != built_clip_tile.key {
+                batches.push(ClipBatch { key: built_clip_tile.key, clips: vec![] });
+            }
+            batches.last_mut().unwrap().clips.push(built_clip_tile.clip);
+        }
+
+        if !batches.is_empty() {
+            self.listener.send(RenderCommand::ClipTiles(batches));
+        }
+    }
+
     fn cull_tiles(&self, paint_metadata: &[PaintMetadata], built_draw_paths: Vec<BuiltDrawPath>)
                   -> CulledTiles {
         let mut culled_tiles = CulledTiles { display_list: vec![] };
@@ -280,7 +310,6 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                                              None,
                                              built_draw_path.blend_mode,
                                              built_draw_path.filter,
-                                             None,
                                              None);
 
                         self.add_alpha_tiles(&mut culled_tiles,
@@ -290,20 +319,7 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                                              color_texture,
                                              built_draw_path.blend_mode,
                                              built_draw_path.filter,
-                                             Some(built_draw_path.mask_0_fill_rule),
+                                             Some(built_draw_path.mask_0_fill_rule));
-                                             None);
-
-                        if let Some(mask_1_fill_rule) = built_draw_path.mask_1_fill_rule {
-                            self.add_alpha_tiles(&mut culled_tiles,
-                                                 layer_z_buffer,
-                                                 &built_draw_path.path.dual_mask_tiles,
-                                                 current_depth,
-                                                 color_texture,
-                                                 built_draw_path.blend_mode,
-                                                 built_draw_path.filter,
-                                                 Some(built_draw_path.mask_0_fill_rule),
-                                                 Some(mask_1_fill_rule));
-                        }
 
                         match built_draw_path.path.solid_tiles {
                             SolidTiles::Regular(ref tiles) => {
@@ -314,7 +330,6 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                                                      color_texture,
                                                      built_draw_path.blend_mode,
                                                      built_draw_path.filter,
-                                                     None,
                                                      None);
                             }
                             SolidTiles::Occluders(_) => {}
@@ -379,8 +394,7 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                        color_texture: Option<TileBatchTexture>,
                        blend_mode: BlendMode,
                        filter: Filter,
-                       mask_0_fill_rule: Option<FillRule>,
+                       mask_0_fill_rule: Option<FillRule>) {
-                       mask_1_fill_rule: Option<FillRule>) {
         let mut batch_indices: Vec<BatchIndex> = vec![];
         for built_alpha_tile in built_alpha_tiles {
             // Early cull if possible.
@@ -406,13 +420,11 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                         blend_mode: batch_blend_mode,
                         filter: batch_filter,
                         mask_0_fill_rule: batch_mask_0_fill_rule,
-                        mask_1_fill_rule: batch_mask_1_fill_rule,
                         tile_page: batch_tile_page
                     })) if *batch_color_texture == color_texture &&
                             batch_blend_mode == blend_mode &&
                             batch_filter == filter &&
                             batch_mask_0_fill_rule == mask_0_fill_rule &&
-                            batch_mask_1_fill_rule == mask_1_fill_rule &&
                             !batch_blend_mode.needs_readable_framebuffer() &&
                             batch_tile_page == built_alpha_tile.page => {
                         dest_batch_index = Some(BatchIndex {
@@ -438,7 +450,6 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                     blend_mode,
                     filter,
                     mask_0_fill_rule,
-                    mask_1_fill_rule,
                     tile_page: built_alpha_tile.page,
                 }));
             }
@@ -504,6 +515,7 @@ impl<'a, 'b> SceneBuilder<'a, 'b> {
                        paint_metadata: &[PaintMetadata],
                        built_draw_paths: Vec<BuiltDrawPath>) {
         self.listener.send(RenderCommand::FlushFills);
+        self.build_clips(&built_draw_paths);
         let culled_tiles = self.cull_tiles(paint_metadata, built_draw_paths);
         self.pack_tiles(culled_tiles);
     }
@@ -567,7 +579,7 @@ impl BuiltPath {
         BuiltPath {
             empty_tiles: vec![],
             single_mask_tiles: vec![],
-            dual_mask_tiles: vec![],
+            clip_tiles: vec![],
             solid_tiles: if occludes {
                 SolidTiles::Occluders(vec![])
             } else {
@@ -688,9 +700,7 @@ impl ObjectBuilder {
             return alpha_tile_id;
         }
 
-        let alpha_tile_index = scene_builder.next_alpha_tile_index.fetch_add(1, Ordering::Relaxed);
+        let alpha_tile_id = AlphaTileId::new(&scene_builder.next_alpha_tile_indices, 0);
-        debug_assert!(alpha_tile_index < u32::MAX as usize);
-        let alpha_tile_id = AlphaTileId(alpha_tile_index as u32);
         self.built_path.tiles.data[local_tile_index].alpha_tile_id = alpha_tile_id;
         alpha_tile_id
     }
@@ -799,31 +809,57 @@ impl ObjectBuilder {
 impl<'a> PackedTile<'a> {
     pub(crate) fn add_to(&self,
                          tiles: &mut Vec<BuiltTile>,
-                         draw_tiling_path_info: &DrawTilingPathInfo) {
+                         clips: &mut Vec<BuiltClip>,
-        let fill_tile_index = self.draw_tile.alpha_tile_id.tile() as u16;
+                         draw_tiling_path_info: &DrawTilingPathInfo,
-        let fill_tile_backdrop = self.draw_tile.backdrop as i8;
+                         scene_builder: &SceneBuilder) {
-        let (clip_tile_index, clip_tile_backdrop) = match self.clip_tile {
+        let draw_tile_page = self.draw_tile.alpha_tile_id.page() as u16;
-            None => (0, 0),
+        let draw_tile_index = self.draw_tile.alpha_tile_id.tile() as u16;
-            Some(clip_tile) => {
+        let draw_tile_backdrop = self.draw_tile.backdrop as i8;
-                // FIXME(pcwalton): This may not always be the case!
-                debug_assert!(!self.draw_tile.alpha_tile_id.is_valid() ||
-                              !clip_tile.alpha_tile_id.is_valid() ||
-                              self.draw_tile.alpha_tile_id.page() ==
-                              clip_tile.alpha_tile_id.page());
 
-                (clip_tile.alpha_tile_id.tile() as u16, clip_tile.backdrop as i8)
+        match self.clip_tile {
+            None => {
+                tiles.push(BuiltTile {
+                    page: draw_tile_page,
+                    tile: Tile::new_alpha(self.tile_coords,
+                                          draw_tile_index,
+                                          draw_tile_backdrop,
+                                          draw_tiling_path_info),
+                });
             }
-        };
+            Some(clip_tile) => {
+                let clip_tile_page = clip_tile.alpha_tile_id.page() as u16;
+                let clip_tile_index = clip_tile.alpha_tile_id.tile() as u16;
+                let clip_tile_backdrop = clip_tile.backdrop;
 
-        tiles.push(BuiltTile {
+                let dest_tile_id = AlphaTileId::new(&scene_builder.next_alpha_tile_indices, 1);
-            page: self.draw_tile.alpha_tile_id.page(),
+                let dest_tile_page = dest_tile_id.page() as u16;
-            tile: Tile::new_alpha(self.tile_coords,
+                let dest_tile_index = dest_tile_id.tile() as u16;
-                                  fill_tile_index,
+
-                                  fill_tile_backdrop,
+                clips.push(BuiltClip {
-                                  clip_tile_index,
+                    clip: Clip::new(dest_tile_index, draw_tile_index, draw_tile_backdrop),
-                                  clip_tile_backdrop,
+                    key: ClipBatchKey {
-                                  draw_tiling_path_info),
+                        src_page: draw_tile_page,
-        });
+                        dest_page: dest_tile_page,
+                        kind: ClipBatchKind::Draw,
+                    },
+                });
+                clips.push(BuiltClip {
+                    clip: Clip::new(dest_tile_index, clip_tile_index, clip_tile_backdrop),
+                    key: ClipBatchKey {
+                        src_page: clip_tile_page,
+                        dest_page: dest_tile_page,
+                        kind: ClipBatchKind::Clip,
+                    },
+                });
+                tiles.push(BuiltTile {
+                    page: dest_tile_page,
+                    tile: Tile::new_alpha(self.tile_coords,
+                                          dest_tile_index,
+                                          0,
+                                          draw_tiling_path_info),
+                });
+            }
+        }
     }
 }
 
@@ -832,21 +868,17 @@ impl Tile {
     fn new_alpha(tile_origin: Vector2I,
                  draw_tile_index: u16,
                  draw_tile_backdrop: i8,
-                 clip_tile_index: u16,
-                 clip_tile_backdrop: i8,
                  draw_tiling_path_info: &DrawTilingPathInfo)
                  -> Tile {
         let mask_0_uv = calculate_mask_uv(draw_tile_index);
-        let mask_1_uv = calculate_mask_uv(clip_tile_index);
         Tile {
             tile_x: tile_origin.x() as i16,
             tile_y: tile_origin.y() as i16,
             mask_0_u: mask_0_uv.x() as u8,
             mask_0_v: mask_0_uv.y() as u8,
-            mask_1_u: mask_1_uv.x() as u8,
-            mask_1_v: mask_1_uv.y() as u8,
             mask_0_backdrop: draw_tile_backdrop,
-            mask_1_backdrop: clip_tile_backdrop,
+            flags: 0,
+            pad: 0,
             color: draw_tiling_path_info.paint_id.0,
         }
     }
@@ -857,6 +889,23 @@ impl Tile {
     }
 }
 
+impl Clip {
+    #[inline]
+    fn new(dest_tile_index: u16, src_tile_index: u16, src_backdrop: i8) -> Clip {
+        let dest_uv = calculate_mask_uv(dest_tile_index);
+        let src_uv = calculate_mask_uv(src_tile_index);
+        Clip {
+            dest_u: dest_uv.x() as u8,
+            dest_v: dest_uv.y() as u8,
+            src_u: src_uv.x() as u8,
+            src_v: src_uv.y() as u8,
+            backdrop: src_backdrop,
+            pad_0: 0,
+            pad_1: 0,
+        }
+    }
+}
+
 fn calculate_mask_uv(tile_index: u16) -> Vector2I {
     debug_assert_eq!(MASK_TILES_ACROSS, MASK_TILES_DOWN);
     let mask_u = tile_index as i32 % MASK_TILES_ACROSS as i32;

renderer/src/gpu/renderer.rs

@@ -10,16 +10,17 @@
 
 use crate::gpu::debug::DebugUIPresenter;
 use crate::gpu::options::{DestFramebuffer, RendererOptions};
-use crate::gpu::shaders::{BlitProgram, BlitVertexArray, CopyTileProgram, CopyTileVertexArray};
+use crate::gpu::shaders::{BlitProgram, BlitVertexArray, ClipTileProgram, ClipTileVertexArray};
-use crate::gpu::shaders::{FillProgram, FillVertexArray, MAX_FILLS_PER_BATCH, ReprojectionProgram};
+use crate::gpu::shaders::{CopyTileProgram, CopyTileVertexArray, FillProgram, FillVertexArray};
-use crate::gpu::shaders::{ReprojectionVertexArray, StencilProgram, StencilVertexArray};
+use crate::gpu::shaders::{MAX_FILLS_PER_BATCH, ReprojectionProgram, ReprojectionVertexArray};
-use crate::gpu::shaders::{TileProgram, TileVertexArray};
+use crate::gpu::shaders::{StencilProgram, StencilVertexArray, TileProgram, TileVertexArray};
-use crate::gpu_data::{Fill, FillBatchEntry, RenderCommand, TextureLocation};
+use crate::gpu_data::{ClipBatch, ClipBatchKey, ClipBatchKind, Fill, FillBatchEntry, RenderCommand};
-use crate::gpu_data::{TextureMetadataEntry, TexturePageDescriptor, TexturePageId};
+use crate::gpu_data::{TextureLocation, TextureMetadataEntry, TexturePageDescriptor, TexturePageId};
 use crate::gpu_data::{Tile, TileBatchTexture};
 use crate::options::BoundingQuad;
 use crate::paint::PaintCompositeOp;
 use crate::tiles::{TILE_HEIGHT, TILE_WIDTH};
+use fxhash::FxHashMap;
 use half::f16;
 use pathfinder_color::{self as color, ColorF, ColorU};
 use pathfinder_content::effects::{BlendMode, BlurDirection, DefringingKernel};
@@ -31,7 +32,7 @@ use pathfinder_geometry::rect::RectI;
 use pathfinder_geometry::transform3d::Transform4F;
 use pathfinder_geometry::util;
 use pathfinder_geometry::vector::{Vector2F, Vector2I, Vector4F, vec2f, vec2i};
-use pathfinder_gpu::{BlendFactor, BlendState, BufferData, BufferTarget, BufferUploadMode};
+use pathfinder_gpu::{BlendFactor, BlendOp, BlendState, BufferData, BufferTarget, BufferUploadMode};
 use pathfinder_gpu::{ClearOps, DepthFunc, DepthState, Device, Primitive, RenderOptions};
 use pathfinder_gpu::{RenderState, RenderTarget, StencilFunc, StencilState, TextureDataRef};
 use pathfinder_gpu::{TextureFormat, UniformData};
@@ -92,7 +93,6 @@ const COMBINER_CTRL_COMPOSITE_COLOR: i32 =          0xe;
 const COMBINER_CTRL_COMPOSITE_LUMINOSITY: i32 =     0xf;
 
 const COMBINER_CTRL_MASK_0_SHIFT: i32 =             0;
-const COMBINER_CTRL_MASK_1_SHIFT: i32 =             2;
 const COMBINER_CTRL_COLOR_FILTER_SHIFT: i32 =       4;
 const COMBINER_CTRL_COLOR_COMBINE_SHIFT: i32 =      6;
 const COMBINER_CTRL_COMPOSITE_SHIFT: i32 =          8;
@@ -111,16 +111,18 @@ where
     fill_program: FillProgram<D>,
     tile_program: TileProgram<D>,
     tile_copy_program: CopyTileProgram<D>,
+    tile_clip_program: ClipTileProgram<D>,
     blit_vertex_array: BlitVertexArray<D>,
     tile_vertex_array: TileVertexArray<D>,
     tile_copy_vertex_array: CopyTileVertexArray<D>,
+    tile_clip_vertex_array: ClipTileVertexArray<D>,
     tile_vertex_buffer: D::Buffer,
     quad_vertex_positions_buffer: D::Buffer,
     quad_vertex_indices_buffer: D::Buffer,
     quads_vertex_indices_buffer: D::Buffer,
     quads_vertex_indices_length: usize,
     fill_vertex_array: FillVertexArray<D>,
-    alpha_tile_pages: Vec<AlphaTilePage<D>>,
+    alpha_tile_pages: FxHashMap<u16, AlphaTilePage<D>>,
     dest_blend_framebuffer: D::Framebuffer,
     intermediate_dest_framebuffer: D::Framebuffer,
     texture_pages: Vec<Option<TexturePage<D>>>,
@@ -167,6 +169,7 @@ where
         let fill_program = FillProgram::new(&device, resources);
         let tile_program = TileProgram::new(&device, resources);
         let tile_copy_program = CopyTileProgram::new(&device, resources);
+        let tile_clip_program = ClipTileProgram::new(&device, resources);
         let stencil_program = StencilProgram::new(&device, resources);
         let reprojection_program = ReprojectionProgram::new(&device, resources);
 
@@ -219,6 +222,12 @@ where
             &tile_vertex_buffer,
             &quads_vertex_indices_buffer,
         );
+        let tile_clip_vertex_array = ClipTileVertexArray::new(
+            &device,
+            &tile_clip_program,
+            &quad_vertex_positions_buffer,
+            &quad_vertex_indices_buffer,
+        );
         let stencil_vertex_array = StencilVertexArray::new(&device, &stencil_program);
         let reprojection_vertex_array = ReprojectionVertexArray::new(
             &device,
@@ -249,16 +258,18 @@ where
             fill_program,
             tile_program,
             tile_copy_program,
+            tile_clip_program,
             blit_vertex_array,
             tile_vertex_array,
             tile_copy_vertex_array,
+            tile_clip_vertex_array,
             tile_vertex_buffer,
             quad_vertex_positions_buffer,
             quad_vertex_indices_buffer,
             quads_vertex_indices_buffer,
             quads_vertex_indices_length: 0,
             fill_vertex_array,
-            alpha_tile_pages: vec![],
+            alpha_tile_pages: FxHashMap::default(),
             dest_blend_framebuffer,
             intermediate_dest_framebuffer,
             texture_pages: vec![],
@@ -291,7 +302,7 @@ where
 
     pub fn begin_scene(&mut self) {
         self.framebuffer_flags = FramebufferFlags::empty();
-        for alpha_tile_page in &mut self.alpha_tile_pages {
+        for alpha_tile_page in self.alpha_tile_pages.values_mut() {
             alpha_tile_page.must_preserve_framebuffer = false;
         }
 
@@ -318,10 +329,14 @@ where
             }
             RenderCommand::AddFills(ref fills) => self.add_fills(fills),
             RenderCommand::FlushFills => {
-                for page_index in 0..(self.alpha_tile_pages.len() as u16) {
+                let page_indices: Vec<_> = self.alpha_tile_pages.keys().cloned().collect();
+                for page_index in page_indices {
                     self.draw_buffered_fills(page_index)
                 }
             }
+            RenderCommand::ClipTiles(ref batches) => {
+                batches.iter().for_each(|batch| self.draw_clip_batch(batch))
+            }
             RenderCommand::BeginTileDrawing => self.begin_tile_drawing(),
             RenderCommand::PushRenderTarget(render_target_id) => {
                 self.push_render_target(render_target_id)
@@ -335,7 +350,6 @@ where
                                 count as u32,
                                 batch.color_texture,
                                 batch.mask_0_fill_rule,
-                                batch.mask_1_fill_rule,
                                 batch.blend_mode,
                                 batch.filter)
             }
@@ -571,20 +585,26 @@ where
 
         for fill_batch_entry in fill_batch {
             let page = fill_batch_entry.page;
-            while self.alpha_tile_pages.len() <= page as usize {
+            if !self.alpha_tile_pages.contains_key(&page) {
-                self.alpha_tile_pages.push(AlphaTilePage::new(&mut self.device));
+                self.alpha_tile_pages.insert(page, AlphaTilePage::new(&mut self.device));
             }
-            self.alpha_tile_pages[page as usize].buffered_fills.push(fill_batch_entry.fill);
+            if self.alpha_tile_pages[&page].buffered_fills.len() == MAX_FILLS_PER_BATCH {
-            if self.alpha_tile_pages[page as usize].buffered_fills.len() == MAX_FILLS_PER_BATCH {
                 self.draw_buffered_fills(page);
             }
+            self.alpha_tile_pages
+                .get_mut(&page)
+                .unwrap()
+                .buffered_fills
+                .push(fill_batch_entry.fill);
         }
     }
 
     fn draw_buffered_fills(&mut self, page: u16) {
         let mask_viewport = self.mask_viewport();
 
-        let alpha_tile_page = &mut self.alpha_tile_pages[page as usize];
+        let alpha_tile_page = self.alpha_tile_pages
+                                  .get_mut(&page)
+                                  .expect("Where's the alpha tile page?");
         let buffered_fills = &mut alpha_tile_page.buffered_fills;
         if buffered_fills.is_empty() {
             return;
@@ -635,6 +655,67 @@ where
         buffered_fills.clear();
     }
 
+    fn draw_clip_batch(&mut self, batch: &ClipBatch) {
+        if batch.clips.is_empty() {
+            return;
+        }
+
+        let ClipBatchKey { dest_page, src_page, kind } = batch.key;
+
+        self.device.allocate_buffer(&self.tile_clip_vertex_array.vertex_buffer,
+                                    BufferData::Memory(&batch.clips),
+                                    BufferTarget::Vertex,
+                                    BufferUploadMode::Dynamic);
+
+        if !self.alpha_tile_pages.contains_key(&dest_page) {
+            self.alpha_tile_pages.insert(dest_page, AlphaTilePage::new(&mut self.device));
+        }
+
+        let mut clear_color = None;
+        if !self.alpha_tile_pages[&dest_page].must_preserve_framebuffer {
+            clear_color = Some(ColorF::default());
+        };
+
+        let blend = match kind {
+            ClipBatchKind::Draw => None,
+            ClipBatchKind::Clip => {
+                Some(BlendState {
+                    src_rgb_factor: BlendFactor::One,
+                    src_alpha_factor: BlendFactor::One,
+                    dest_rgb_factor: BlendFactor::One,
+                    dest_alpha_factor: BlendFactor::One,
+                    op: BlendOp::Min,
+                })
+            }
+        };
+
+        let mask_viewport = self.mask_viewport();
+
+        {
+            let dest_framebuffer = &self.alpha_tile_pages[&dest_page].framebuffer;
+            let src_framebuffer = &self.alpha_tile_pages[&src_page].framebuffer;
+            let src_texture = self.device.framebuffer_texture(&src_framebuffer);
+
+            debug_assert!(batch.clips.len() <= u32::MAX as usize);
+            self.device.draw_elements_instanced(6, batch.clips.len() as u32, &RenderState {
+                target: &RenderTarget::Framebuffer(dest_framebuffer),
+                program: &self.tile_clip_program.program,
+                vertex_array: &self.tile_clip_vertex_array.vertex_array,
+                primitive: Primitive::Triangles,
+                textures: &[src_texture],
+                uniforms: &[(&self.tile_clip_program.src_uniform, UniformData::TextureUnit(0))],
+                viewport: mask_viewport,
+                options: RenderOptions {
+                    blend,
+                    clear_ops: ClearOps { color: clear_color, ..ClearOps::default() },
+                    ..RenderOptions::default()
+                },
+            });
+        }
+
+        self.alpha_tile_pages.get_mut(&dest_page).unwrap().must_preserve_framebuffer = true;
+    }
+
     fn tile_transform(&self) -> Transform4F {
         let draw_viewport = self.draw_viewport().size().to_f32();
         let scale = Vector4F::new(2.0 / draw_viewport.x(), -2.0 / draw_viewport.y(), 1.0, 1.0);
@@ -646,7 +727,6 @@ where
                   tile_count: u32,
                   color_texture_0: Option<TileBatchTexture>,
                   mask_0_fill_rule: Option<FillRule>,
-                  mask_1_fill_rule: Option<FillRule>,
                   blend_mode: BlendMode,
                   filter: Filter) {
         // TODO(pcwalton): Disable blend for solid tiles.
@@ -660,14 +740,13 @@ where
         let draw_viewport = self.draw_viewport();
 
         let mut ctrl = 0;
-        for &(fill_rule, shift) in &[
+        match mask_0_fill_rule {
-            (mask_0_fill_rule, COMBINER_CTRL_MASK_0_SHIFT),
+            None => {}
-            (mask_1_fill_rule, COMBINER_CTRL_MASK_1_SHIFT),
+            Some(FillRule::Winding) => {
-        ] {
+                ctrl |= COMBINER_CTRL_MASK_WINDING << COMBINER_CTRL_MASK_0_SHIFT
-            match fill_rule {
+            }
-                None => {}
+            Some(FillRule::EvenOdd) => {
-                Some(FillRule::Winding) => ctrl |= COMBINER_CTRL_MASK_WINDING << shift,
+                ctrl |= COMBINER_CTRL_MASK_EVEN_ODD << COMBINER_CTRL_MASK_0_SHIFT
-                Some(FillRule::EvenOdd) => ctrl |= COMBINER_CTRL_MASK_EVEN_ODD << shift,
             }
         }
 
@@ -695,13 +774,7 @@ where
             uniforms.push((&self.tile_program.mask_texture_0_uniform,
                            UniformData::TextureUnit(textures.len() as u32)));
             textures.push(self.device.framebuffer_texture(
-                &self.alpha_tile_pages[tile_page as usize].framebuffer));
+                &self.alpha_tile_pages[&tile_page].framebuffer));
-        }
-        if mask_1_fill_rule.is_some() {
-            uniforms.push((&self.tile_program.mask_texture_1_uniform,
-                           UniformData::TextureUnit(textures.len() as u32)));
-            textures.push(self.device.framebuffer_texture(
-                &self.alpha_tile_pages[tile_page as usize].framebuffer));
         }
 
         // TODO(pcwalton): Refactor.

renderer/src/gpu/shaders.rs

@@ -16,6 +16,7 @@ use pathfinder_resources::ResourceLoader;
 // TODO(pcwalton): Replace with `mem::size_of` calls?
 const FILL_INSTANCE_SIZE: usize = 8;
 const TILE_INSTANCE_SIZE: usize = 12;
+const CLIP_TILE_INSTANCE_SIZE: usize = 8;
 
 pub const MAX_FILLS_PER_BATCH: usize = 0x4000;
 
@@ -165,8 +166,6 @@ impl<D> TileVertexArray<D> where D: Device {
             device.get_vertex_attr(&tile_program.program, "TileOrigin").unwrap();
         let mask_0_tex_coord_attr =
             device.get_vertex_attr(&tile_program.program, "MaskTexCoord0").unwrap();
-        let mask_1_tex_coord_attr =
-            device.get_vertex_attr(&tile_program.program, "MaskTexCoord1").unwrap();
         let mask_backdrop_attr =
             device.get_vertex_attr(&tile_program.program, "MaskBackdrop").unwrap();
         let color_attr = device.get_vertex_attr(&tile_program.program, "Color").unwrap();
@@ -200,21 +199,12 @@ impl<D> TileVertexArray<D> where D: Device {
             divisor: 1,
             buffer_index: 1,
         });
-        device.configure_vertex_attr(&vertex_array, &mask_1_tex_coord_attr, &VertexAttrDescriptor {
-            size: 2,
-            class: VertexAttrClass::Int,
-            attr_type: VertexAttrType::U8,
-            stride: TILE_INSTANCE_SIZE,
-            offset: 6,
-            divisor: 1,
-            buffer_index: 1,
-        });
         device.configure_vertex_attr(&vertex_array, &mask_backdrop_attr, &VertexAttrDescriptor {
             size: 2,
             class: VertexAttrClass::Int,
             attr_type: VertexAttrType::I8,
             stride: TILE_INSTANCE_SIZE,
-            offset: 8,
+            offset: 6,
             divisor: 1,
             buffer_index: 1,
         });
@@ -223,7 +213,7 @@ impl<D> TileVertexArray<D> where D: Device {
             class: VertexAttrClass::Int,
             attr_type: VertexAttrType::I16,
             stride: TILE_INSTANCE_SIZE,
-            offset: 10,
+            offset: 8,
             divisor: 1,
             buffer_index: 1,
         });
@@ -265,6 +255,74 @@ impl<D> CopyTileVertexArray<D> where D: Device {
     }
 }
 
+pub struct ClipTileVertexArray<D> where D: Device {
+    pub vertex_array: D::VertexArray,
+    pub vertex_buffer: D::Buffer,
+}
+
+impl<D> ClipTileVertexArray<D> where D: Device {
+    pub fn new(device: &D,
+               clip_tile_program: &ClipTileProgram<D>,
+               quad_vertex_positions_buffer: &D::Buffer,
+               quad_vertex_indices_buffer: &D::Buffer)
+               -> ClipTileVertexArray<D> {
+        let vertex_array = device.create_vertex_array();
+        let vertex_buffer = device.create_buffer();
+
+        let tile_offset_attr =
+            device.get_vertex_attr(&clip_tile_program.program, "TileOffset").unwrap();
+        let dest_tile_origin_attr =
+            device.get_vertex_attr(&clip_tile_program.program, "DestTileOrigin").unwrap();
+        let src_tile_origin_attr =
+            device.get_vertex_attr(&clip_tile_program.program, "SrcTileOrigin").unwrap();
+        let src_backdrop_attr =
+            device.get_vertex_attr(&clip_tile_program.program, "SrcBackdrop").unwrap();
+
+        device.bind_buffer(&vertex_array, quad_vertex_positions_buffer, BufferTarget::Vertex);
+        device.configure_vertex_attr(&vertex_array, &tile_offset_attr, &VertexAttrDescriptor {
+            size: 2,
+            class: VertexAttrClass::Int,
+            attr_type: VertexAttrType::I16,
+            stride: 4,
+            offset: 0,
+            divisor: 0,
+            buffer_index: 0,
+        });
+        device.bind_buffer(&vertex_array, &vertex_buffer, BufferTarget::Vertex);
+        device.configure_vertex_attr(&vertex_array, &dest_tile_origin_attr, &VertexAttrDescriptor {
+            size: 2,
+            class: VertexAttrClass::Int,
+            attr_type: VertexAttrType::U8,
+            stride: CLIP_TILE_INSTANCE_SIZE,
+            offset: 0,
+            divisor: 1,
+            buffer_index: 1,
+        });
+        device.configure_vertex_attr(&vertex_array, &src_tile_origin_attr, &VertexAttrDescriptor {
+            size: 2,
+            class: VertexAttrClass::Int,
+            attr_type: VertexAttrType::U8,
+            stride: CLIP_TILE_INSTANCE_SIZE,
+            offset: 2,
+            divisor: 1,
+            buffer_index: 1,
+        });
+        device.configure_vertex_attr(&vertex_array, &src_backdrop_attr, &VertexAttrDescriptor {
+            size: 1,
+            class: VertexAttrClass::Int,
+            attr_type: VertexAttrType::I8,
+            stride: CLIP_TILE_INSTANCE_SIZE,
+            offset: 4,
+            divisor: 1,
+            buffer_index: 1,
+        });
+        device.bind_buffer(&vertex_array, quad_vertex_indices_buffer, BufferTarget::Index);
+
+        ClipTileVertexArray { vertex_array, vertex_buffer }
+    }
+}
+
+
 pub struct BlitProgram<D> where D: Device {
     pub program: D::Program,
     pub src_uniform: D::Uniform,
@@ -316,7 +374,6 @@ pub struct TileProgram<D> where D: Device {
     pub color_texture_0_uniform: D::Uniform,
     pub color_texture_1_uniform: D::Uniform,
     pub mask_texture_0_uniform: D::Uniform,
-    pub mask_texture_1_uniform: D::Uniform,
     pub gamma_lut_uniform: D::Uniform,
     pub color_texture_0_size_uniform: D::Uniform,
     pub filter_params_0_uniform: D::Uniform,
@@ -337,7 +394,6 @@ impl<D> TileProgram<D> where D: Device {
         let color_texture_0_uniform = device.get_uniform(&program, "ColorTexture0");
         let color_texture_1_uniform = device.get_uniform(&program, "ColorTexture1");
         let mask_texture_0_uniform = device.get_uniform(&program, "MaskTexture0");
-        let mask_texture_1_uniform = device.get_uniform(&program, "MaskTexture1");
         let gamma_lut_uniform = device.get_uniform(&program, "GammaLUT");
         let color_texture_0_size_uniform = device.get_uniform(&program, "ColorTexture0Size");
         let filter_params_0_uniform = device.get_uniform(&program, "FilterParams0");
@@ -355,7 +411,6 @@ impl<D> TileProgram<D> where D: Device {
             color_texture_0_uniform,
             color_texture_1_uniform,
             mask_texture_0_uniform,
-            mask_texture_1_uniform,
             gamma_lut_uniform,
             color_texture_0_size_uniform,
             filter_params_0_uniform,
@@ -392,6 +447,19 @@ impl<D> CopyTileProgram<D> where D: Device {
     }
 }
 
+pub struct ClipTileProgram<D> where D: Device {
+    pub program: D::Program,
+    pub src_uniform: D::Uniform,
+}
+
+impl<D> ClipTileProgram<D> where D: Device {
+    pub fn new(device: &D, resources: &dyn ResourceLoader) -> ClipTileProgram<D> {
+        let program = device.create_program(resources, "tile_clip");
+        let src_uniform = device.get_uniform(&program, "Src");
+        ClipTileProgram { program, src_uniform }
+    }
+}
+
 pub struct StencilProgram<D>
 where
     D: Device,

renderer/src/gpu_data.rs

@@ -10,6 +10,7 @@
 
 //! Packed data ready to be sent to the GPU.
 
+use crate::builder::{ALPHA_TILES_PER_LEVEL, ALPHA_TILE_LEVEL_COUNT};
 use crate::options::BoundingQuad;
 use crate::paint::PaintCompositeOp;
 use pathfinder_color::ColorU;
@@ -23,7 +24,9 @@ use pathfinder_geometry::vector::Vector2I;
 use pathfinder_gpu::TextureSamplingFlags;
 use std::fmt::{Debug, Formatter, Result as DebugResult};
 use std::sync::Arc;
+use std::sync::atomic::{AtomicUsize, Ordering};
 use std::time::Duration;
+use std::u32;
 
 pub enum RenderCommand {
     // Starts rendering a frame.
@@ -61,6 +64,9 @@ pub enum RenderCommand {
     // Flushes the queue of fills.
     FlushFills,
 
+    // Renders clips to the mask tile.
+    ClipTiles(Vec<ClipBatch>),
+
     // Pushes a render target onto the stack. Draw commands go to the render target on top of the
     // stack.
     PushRenderTarget(RenderTargetId),
@@ -97,7 +103,6 @@ pub struct TileBatch {
     pub tiles: Vec<Tile>,
     pub color_texture: Option<TileBatchTexture>,
     pub mask_0_fill_rule: Option<FillRule>,
-    pub mask_1_fill_rule: Option<FillRule>,
     pub filter: Filter,
     pub blend_mode: BlendMode,
     pub tile_page: u16,
@@ -147,6 +152,38 @@ pub struct Fill {
     pub alpha_tile_index: u16,
 }
 
+#[derive(Clone, Debug)]
+pub struct ClipBatch {
+    pub clips: Vec<Clip>,
+    pub key: ClipBatchKey,
+}
+
+#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord)]
+pub struct ClipBatchKey {
+    pub dest_page: u16,
+    pub src_page: u16,
+    pub kind: ClipBatchKind,
+}
+
+// Order is significant here.
+#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord)]
+pub enum ClipBatchKind {
+    Draw,
+    Clip,
+}
+
+#[derive(Clone, Copy, Debug, Default)]
+#[repr(C)]
+pub struct Clip {
+    pub dest_u: u8,
+    pub dest_v: u8,
+    pub src_u: u8,
+    pub src_v: u8,
+    pub backdrop: i8,
+    pub pad_0: u8,
+    pub pad_1: u16,
+}
+
 #[derive(Clone, Copy, Debug, Default)]
 #[repr(C)]
 pub struct Tile {
@@ -154,17 +191,24 @@ pub struct Tile {
     pub tile_y: i16,
     pub mask_0_u: u8,
     pub mask_0_v: u8,
-    pub mask_1_u: u8,
-    pub mask_1_v: u8,
     pub mask_0_backdrop: i8,
-    pub mask_1_backdrop: i8,
+    pub pad: u8,
     pub color: u16,
+    pub flags: u16,
 }
 
 #[derive(Clone, Copy, PartialEq, Debug)]
 pub struct AlphaTileId(pub u32);
 
 impl AlphaTileId {
+    #[inline]
+    pub fn new(next_alpha_tile_index: &[AtomicUsize; ALPHA_TILE_LEVEL_COUNT], level: usize) 
+               -> AlphaTileId {
+        let alpha_tile_index = next_alpha_tile_index[level].fetch_add(1, Ordering::Relaxed);
+        debug_assert!(alpha_tile_index < ALPHA_TILES_PER_LEVEL);
+        AlphaTileId((level * ALPHA_TILES_PER_LEVEL + alpha_tile_index) as u32)
+    }
+
     #[inline]
     pub fn invalid() -> AlphaTileId {
         AlphaTileId(!0)
@@ -206,6 +250,9 @@ impl Debug for RenderCommand {
                 write!(formatter, "AddFills(x{})", fills.len())
             }
             RenderCommand::FlushFills => write!(formatter, "FlushFills"),
+            RenderCommand::ClipTiles(ref batches) => {
+                write!(formatter, "ClipTiles(x{})", batches.len())
+            }
             RenderCommand::PushRenderTarget(render_target_id) => {
                 write!(formatter, "PushRenderTarget({:?})", render_target_id)
             }

renderer/src/tiles.rs

@@ -146,23 +146,25 @@ impl<'a, 'b> Tiler<'a, 'b> {
                             occluders.push(Occluder::new(packed_tile.tile_coords));
                         }
                         SolidTiles::Regular(ref mut solid_tiles) => {
-                            packed_tile.add_to(solid_tiles, &draw_tiling_path_info);
+                            packed_tile.add_to(solid_tiles,
+                                               &mut self.object_builder.built_path.clip_tiles,
+                                               &draw_tiling_path_info,
+                                               &self.scene_builder);
                         }
                     }
                 }
                 TileType::SingleMask => {
                     debug_assert_ne!(packed_tile.draw_tile.alpha_tile_id.page(), !0);
                     packed_tile.add_to(&mut self.object_builder.built_path.single_mask_tiles,
-                                       &draw_tiling_path_info);
+                                       &mut self.object_builder.built_path.clip_tiles,
-                }
+                                       &draw_tiling_path_info,
-                TileType::DualMask => {
+                                       &self.scene_builder);
-                    debug_assert_ne!(packed_tile.draw_tile.alpha_tile_id.page(), !0);
-                    packed_tile.add_to(&mut self.object_builder.built_path.dual_mask_tiles,
-                                       &draw_tiling_path_info);
                 }
                 TileType::Empty if blend_mode_is_destructive => {
                     packed_tile.add_to(&mut self.object_builder.built_path.empty_tiles,
-                                       &draw_tiling_path_info);
+                                       &mut self.object_builder.built_path.clip_tiles,
+                                       &draw_tiling_path_info,
+                                       &self.scene_builder);
                 }
                 TileType::Empty => {
                     // Just cull.
@@ -404,7 +406,6 @@ pub(crate) enum TileType {
     Solid,
     Empty,
     SingleMask,
-    DualMask,
 }
 
 impl<'a> PackedTile<'a> {
@@ -504,7 +505,7 @@ impl<'a> PackedTile<'a> {
             Some(clip_tile) => {
                 // We have both a draw and clip mask. Composite them together.
                 PackedTile {
-                    tile_type: TileType::DualMask,
+                    tile_type: TileType::SingleMask,
                     tile_coords,
                     draw_tile,
                     clip_tile: Some(clip_tile),

renderer/src/z_buffer.rs

@@ -92,7 +92,6 @@ impl ZBuffer {
                         filter: Filter::None,
                         blend_mode: BlendMode::default(),
                         mask_0_fill_rule: None,
-                        mask_1_fill_rule: None,
                         tile_page: !0,
                     });
                 }
@@ -112,11 +111,10 @@ impl Tile {
             tile_x: tile_origin.x() as i16,
             tile_y: tile_origin.y() as i16,
             mask_0_backdrop: 0,
-            mask_1_backdrop: 0,
             mask_0_u: 0,
             mask_0_v: 0,
-            mask_1_u: 0,
+            flags: 0,
-            mask_1_v: 0,
+            pad: 0,
             color: paint_id.0,
         }
     }

resources/shaders/gl3/tile.fs.glsl

@@ -77,12 +77,10 @@ precision highp sampler2D;
 
 
 
-
 
 
 uniform sampler2D uColorTexture0;
 uniform sampler2D uMaskTexture0;
-uniform sampler2D uMaskTexture1;
 uniform sampler2D uDestTexture;
 uniform sampler2D uGammaLUT;
 uniform vec4 uFilterParams0;
@@ -93,7 +91,6 @@ uniform vec2 uColorTexture0Size;
 uniform int uCtrl;
 
 in vec3 vMaskTexCoord0;
-in vec3 vMaskTexCoord1;
 in vec2 vColorTexCoord0;
 in vec4 vBaseColor;
 
@@ -562,10 +559,8 @@ float sampleMask(float maskAlpha,
 void calculateColor(int ctrl){
 
     int maskCtrl0 =(ctrl >> 0)& 0x3;
-    int maskCtrl1 =(ctrl >> 2)& 0x3;
     float maskAlpha = 1.0;
     maskAlpha = sampleMask(maskAlpha, uMaskTexture0, vMaskTexCoord0, maskCtrl0);
-    maskAlpha = sampleMask(maskAlpha, uMaskTexture1, vMaskTexCoord1, maskCtrl1);
 
 
     vec4 color = vBaseColor;

resources/shaders/gl3/tile.vs.glsl

@@ -23,12 +23,10 @@ uniform ivec2 uTextureMetadataSize;
 in ivec2 aTileOffset;
 in ivec2 aTileOrigin;
 in uvec2 aMaskTexCoord0;
-in uvec2 aMaskTexCoord1;
 in ivec2 aMaskBackdrop;
 in int aColor;
 
 out vec3 vMaskTexCoord0;
-out vec3 vMaskTexCoord1;
 out vec2 vColorTexCoord0;
 out vec4 vBaseColor;
 
@@ -37,7 +35,6 @@ void main(){
     vec2 position =(tileOrigin + tileOffset)* uTileSize;
 
     vec2 maskTexCoord0 =(vec2(aMaskTexCoord0)+ tileOffset)/ 256.0;
-    vec2 maskTexCoord1 =(vec2(aMaskTexCoord1)+ tileOffset)/ 256.0;
 
     vec2 textureMetadataScale = vec2(1.0)/ vec2(uTextureMetadataSize);
     vec2 metadataEntryCoord = vec2(aColor % 128 * 4, aColor / 128);
@@ -50,7 +47,6 @@ void main(){
 
     vColorTexCoord0 = mat2(colorTexMatrix0)* position + colorTexOffsets . xy;
     vMaskTexCoord0 = vec3(maskTexCoord0, float(aMaskBackdrop . x));
-    vMaskTexCoord1 = vec3(maskTexCoord1, float(aMaskBackdrop . y));
     vBaseColor = baseColor;
     gl_Position = uTransform * vec4(position, 0.0, 1.0);
 }

resources/shaders/metal/tile.fs.metal

@@ -10,20 +10,18 @@ struct spvDescriptorSetBuffer0
 {
     texture2d<float> uMaskTexture0 [[id(0)]];
     sampler uMaskTexture0Smplr [[id(1)]];
-    texture2d<float> uMaskTexture1 [[id(2)]];
+    texture2d<float> uColorTexture0 [[id(2)]];
-    sampler uMaskTexture1Smplr [[id(3)]];
+    sampler uColorTexture0Smplr [[id(3)]];
-    texture2d<float> uColorTexture0 [[id(4)]];
+    texture2d<float> uGammaLUT [[id(4)]];
-    sampler uColorTexture0Smplr [[id(5)]];
+    sampler uGammaLUTSmplr [[id(5)]];
-    texture2d<float> uGammaLUT [[id(6)]];
+    constant float2* uColorTexture0Size [[id(6)]];
-    sampler uGammaLUTSmplr [[id(7)]];
+    constant float2* uFramebufferSize [[id(7)]];
-    constant float2* uColorTexture0Size [[id(8)]];
+    constant float4* uFilterParams0 [[id(8)]];
-    constant float2* uFramebufferSize [[id(9)]];
+    constant float4* uFilterParams1 [[id(9)]];
-    constant float4* uFilterParams0 [[id(10)]];
+    constant float4* uFilterParams2 [[id(10)]];
-    constant float4* uFilterParams1 [[id(11)]];
+    texture2d<float> uDestTexture [[id(11)]];
-    constant float4* uFilterParams2 [[id(12)]];
+    sampler uDestTextureSmplr [[id(12)]];
-    texture2d<float> uDestTexture [[id(13)]];
+    constant int* uCtrl [[id(13)]];
-    sampler uDestTextureSmplr [[id(14)]];
-    constant int* uCtrl [[id(15)]];
 };
 
 constant float3 _1040 = {};
@@ -36,9 +34,8 @@ struct main0_out
 struct main0_in
 {
     float3 vMaskTexCoord0 [[user(locn0)]];
-    float3 vMaskTexCoord1 [[user(locn1)]];
+    float2 vColorTexCoord0 [[user(locn1)]];
-    float2 vColorTexCoord0 [[user(locn2)]];
+    float4 vBaseColor [[user(locn2)]];
-    float4 vBaseColor [[user(locn3)]];
 };
 
 // Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
@@ -548,47 +545,42 @@ float4 composite(thread const float4& srcColor, thread const texture2d<float> de
     return float4(((srcColor.xyz * (srcColor.w * (1.0 - destColor.w))) + (blendedRGB * (srcColor.w * destColor.w))) + (destColor.xyz * (1.0 - srcColor.w)), 1.0);
 }
 
-void calculateColor(thread const int& ctrl, thread texture2d<float> uMaskTexture0, thread const sampler uMaskTexture0Smplr, thread float3& vMaskTexCoord0, thread texture2d<float> uMaskTexture1, thread const sampler uMaskTexture1Smplr, thread float3& vMaskTexCoord1, thread float4& vBaseColor, thread float2& vColorTexCoord0, thread texture2d<float> uColorTexture0, thread const sampler uColorTexture0Smplr, thread texture2d<float> uGammaLUT, thread const sampler uGammaLUTSmplr, thread float2 uColorTexture0Size, thread float4& gl_FragCoord, thread float2 uFramebufferSize, thread float4 uFilterParams0, thread float4 uFilterParams1, thread float4 uFilterParams2, thread texture2d<float> uDestTexture, thread const sampler uDestTextureSmplr, thread float4& oFragColor)
+void calculateColor(thread const int& ctrl, thread texture2d<float> uMaskTexture0, thread const sampler uMaskTexture0Smplr, thread float3& vMaskTexCoord0, thread float4& vBaseColor, thread float2& vColorTexCoord0, thread texture2d<float> uColorTexture0, thread const sampler uColorTexture0Smplr, thread texture2d<float> uGammaLUT, thread const sampler uGammaLUTSmplr, thread float2 uColorTexture0Size, thread float4& gl_FragCoord, thread float2 uFramebufferSize, thread float4 uFilterParams0, thread float4 uFilterParams1, thread float4 uFilterParams2, thread texture2d<float> uDestTexture, thread const sampler uDestTextureSmplr, thread float4& oFragColor)
 {
     int maskCtrl0 = (ctrl >> 0) & 3;
-    int maskCtrl1 = (ctrl >> 2) & 3;
     float maskAlpha = 1.0;
     float param = maskAlpha;
     float3 param_1 = vMaskTexCoord0;
     int param_2 = maskCtrl0;
     maskAlpha = sampleMask(param, uMaskTexture0, uMaskTexture0Smplr, param_1, param_2);
-    float param_3 = maskAlpha;
-    float3 param_4 = vMaskTexCoord1;
-    int param_5 = maskCtrl1;
-    maskAlpha = sampleMask(param_3, uMaskTexture1, uMaskTexture1Smplr, param_4, param_5);
     float4 color = vBaseColor;
     int color0Combine = (ctrl >> 6) & 3;
     if (color0Combine != 0)
     {
         int color0Filter = (ctrl >> 4) & 3;
-        float2 param_6 = vColorTexCoord0;
+        float2 param_3 = vColorTexCoord0;
-        float2 param_7 = uColorTexture0Size;
+        float2 param_4 = uColorTexture0Size;
-        float2 param_8 = gl_FragCoord.xy;
+        float2 param_5 = gl_FragCoord.xy;
-        float2 param_9 = uFramebufferSize;
+        float2 param_6 = uFramebufferSize;
-        float4 param_10 = uFilterParams0;
+        float4 param_7 = uFilterParams0;
-        float4 param_11 = uFilterParams1;
+        float4 param_8 = uFilterParams1;
-        float4 param_12 = uFilterParams2;
+        float4 param_9 = uFilterParams2;
-        int param_13 = color0Filter;
+        int param_10 = color0Filter;
-        float4 color0 = filterColor(param_6, uColorTexture0, uColorTexture0Smplr, uGammaLUT, uGammaLUTSmplr, param_7, param_8, param_9, param_10, param_11, param_12, param_13);
+        float4 color0 = filterColor(param_3, uColorTexture0, uColorTexture0Smplr, uGammaLUT, uGammaLUTSmplr, param_4, param_5, param_6, param_7, param_8, param_9, param_10);
-        float4 param_14 = color;
+        float4 param_11 = color;
-        float4 param_15 = color0;
+        float4 param_12 = color0;
-        int param_16 = color0Combine;
+        int param_13 = color0Combine;
-        color = combineColor0(param_14, param_15, param_16);
+        color = combineColor0(param_11, param_12, param_13);
     }
     color.w *= maskAlpha;
     int compositeOp = (ctrl >> 8) & 15;
-    float4 param_17 = color;
+    float4 param_14 = color;
-    float2 param_18 = uFramebufferSize;
+    float2 param_15 = uFramebufferSize;
-    float2 param_19 = gl_FragCoord.xy;
+    float2 param_16 = gl_FragCoord.xy;
-    int param_20 = compositeOp;
+    int param_17 = compositeOp;
-    color = composite(param_17, uDestTexture, uDestTextureSmplr, param_18, param_19, param_20);
+    color = composite(param_14, uDestTexture, uDestTextureSmplr, param_15, param_16, param_17);
-    float3 _1337 = color.xyz * color.w;
+    float3 _1324 = color.xyz * color.w;
-    color = float4(_1337.x, _1337.y, _1337.z, color.w);
+    color = float4(_1324.x, _1324.y, _1324.z, color.w);
     oFragColor = color;
 }
 
@@ -596,7 +588,7 @@ fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuff
 {
     main0_out out = {};
     int param = (*spvDescriptorSet0.uCtrl);
-    calculateColor(param, spvDescriptorSet0.uMaskTexture0, spvDescriptorSet0.uMaskTexture0Smplr, in.vMaskTexCoord0, spvDescriptorSet0.uMaskTexture1, spvDescriptorSet0.uMaskTexture1Smplr, in.vMaskTexCoord1, in.vBaseColor, in.vColorTexCoord0, spvDescriptorSet0.uColorTexture0, spvDescriptorSet0.uColorTexture0Smplr, spvDescriptorSet0.uGammaLUT, spvDescriptorSet0.uGammaLUTSmplr, (*spvDescriptorSet0.uColorTexture0Size), gl_FragCoord, (*spvDescriptorSet0.uFramebufferSize), (*spvDescriptorSet0.uFilterParams0), (*spvDescriptorSet0.uFilterParams1), (*spvDescriptorSet0.uFilterParams2), spvDescriptorSet0.uDestTexture, spvDescriptorSet0.uDestTextureSmplr, out.oFragColor);
+    calculateColor(param, spvDescriptorSet0.uMaskTexture0, spvDescriptorSet0.uMaskTexture0Smplr, in.vMaskTexCoord0, in.vBaseColor, in.vColorTexCoord0, spvDescriptorSet0.uColorTexture0, spvDescriptorSet0.uColorTexture0Smplr, spvDescriptorSet0.uGammaLUT, spvDescriptorSet0.uGammaLUTSmplr, (*spvDescriptorSet0.uColorTexture0Size), gl_FragCoord, (*spvDescriptorSet0.uFramebufferSize), (*spvDescriptorSet0.uFilterParams0), (*spvDescriptorSet0.uFilterParams1), (*spvDescriptorSet0.uFilterParams2), spvDescriptorSet0.uDestTexture, spvDescriptorSet0.uDestTextureSmplr, out.oFragColor);
     return out;
 }
 

resources/shaders/metal/tile.vs.metal

@@ -16,9 +16,8 @@ struct spvDescriptorSetBuffer0
 struct main0_out
 {
     float3 vMaskTexCoord0 [[user(locn0)]];
-    float3 vMaskTexCoord1 [[user(locn1)]];
+    float2 vColorTexCoord0 [[user(locn1)]];
-    float2 vColorTexCoord0 [[user(locn2)]];
+    float4 vBaseColor [[user(locn2)]];
-    float4 vBaseColor [[user(locn3)]];
     float4 gl_Position [[position]];
 };
 
@@ -27,9 +26,8 @@ struct main0_in
     int2 aTileOffset [[attribute(0)]];
     int2 aTileOrigin [[attribute(1)]];
     uint2 aMaskTexCoord0 [[attribute(2)]];
-    uint2 aMaskTexCoord1 [[attribute(3)]];
+    int2 aMaskBackdrop [[attribute(3)]];
-    int2 aMaskBackdrop [[attribute(4)]];
+    int aColor [[attribute(4)]];
-    int aColor [[attribute(5)]];
 };
 
 vertex main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]])
@@ -39,7 +37,6 @@ vertex main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer
     float2 tileOffset = float2(in.aTileOffset);
     float2 position = (tileOrigin + tileOffset) * (*spvDescriptorSet0.uTileSize);
     float2 maskTexCoord0 = (float2(in.aMaskTexCoord0) + tileOffset) / float2(256.0);
-    float2 maskTexCoord1 = (float2(in.aMaskTexCoord1) + tileOffset) / float2(256.0);
     float2 textureMetadataScale = float2(1.0) / float2((*spvDescriptorSet0.uTextureMetadataSize));
     float2 metadataEntryCoord = float2(float((in.aColor % 128) * 4), float(in.aColor / 128));
     float2 colorTexMatrix0Coord = (metadataEntryCoord + float2(0.5)) * textureMetadataScale;
@@ -50,7 +47,6 @@ vertex main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer
     float4 baseColor = spvDescriptorSet0.uTextureMetadata.sample(spvDescriptorSet0.uTextureMetadataSmplr, baseColorCoord, level(0.0));
     out.vColorTexCoord0 = (float2x2(float2(colorTexMatrix0.xy), float2(colorTexMatrix0.zw)) * position) + colorTexOffsets.xy;
     out.vMaskTexCoord0 = float3(maskTexCoord0, float(in.aMaskBackdrop.x));
-    out.vMaskTexCoord1 = float3(maskTexCoord1, float(in.aMaskBackdrop.y));
     out.vBaseColor = baseColor;
     out.gl_Position = (*spvDescriptorSet0.uTransform) * float4(position, 0.0, 1.0);
     return out;

shaders/Makefile

@@ -19,6 +19,8 @@ SHADERS=\
 	stencil.vs.glsl \
 	tile.fs.glsl \
 	tile.vs.glsl \
+	tile_clip.fs.glsl \
+	tile_clip.vs.glsl \
 	tile_copy.fs.glsl \
 	tile_copy.vs.glsl \
 	$(EMPTY)

shaders/tile.fs.glsl

@@ -73,14 +73,12 @@ precision highp sampler2D;
 #define COMBINER_CTRL_COMPOSITE_LUMINOSITY      0xf
 
 #define COMBINER_CTRL_MASK_0_SHIFT              0
-#define COMBINER_CTRL_MASK_1_SHIFT              2
 #define COMBINER_CTRL_COLOR_FILTER_SHIFT        4
 #define COMBINER_CTRL_COLOR_COMBINE_SHIFT       6
 #define COMBINER_CTRL_COMPOSITE_SHIFT           8
 
 uniform sampler2D uColorTexture0;
 uniform sampler2D uMaskTexture0;
-uniform sampler2D uMaskTexture1;
 uniform sampler2D uDestTexture;
 uniform sampler2D uGammaLUT;
 uniform vec4 uFilterParams0;
@@ -91,7 +89,6 @@ uniform vec2 uColorTexture0Size;
 uniform int uCtrl;
 
 in vec3 vMaskTexCoord0;
-in vec3 vMaskTexCoord1;
 in vec2 vColorTexCoord0;
 in vec4 vBaseColor;
 
@@ -560,10 +557,8 @@ float sampleMask(float maskAlpha,
 void calculateColor(int ctrl) {
     // Sample mask.
     int maskCtrl0 = (ctrl >> COMBINER_CTRL_MASK_0_SHIFT) & COMBINER_CTRL_MASK_MASK;
-    int maskCtrl1 = (ctrl >> COMBINER_CTRL_MASK_1_SHIFT) & COMBINER_CTRL_MASK_MASK;
     float maskAlpha = 1.0;
     maskAlpha = sampleMask(maskAlpha, uMaskTexture0, vMaskTexCoord0, maskCtrl0);
-    maskAlpha = sampleMask(maskAlpha, uMaskTexture1, vMaskTexCoord1, maskCtrl1);
 
     // Sample color.
     vec4 color = vBaseColor;

shaders/tile.vs.glsl

@@ -21,12 +21,10 @@ uniform ivec2 uTextureMetadataSize;
 in ivec2 aTileOffset;
 in ivec2 aTileOrigin;
 in uvec2 aMaskTexCoord0;
-in uvec2 aMaskTexCoord1;
 in ivec2 aMaskBackdrop;
 in int aColor;
 
 out vec3 vMaskTexCoord0;
-out vec3 vMaskTexCoord1;
 out vec2 vColorTexCoord0;
 out vec4 vBaseColor;
 
@@ -35,7 +33,6 @@ void main() {
     vec2 position = (tileOrigin + tileOffset) * uTileSize;
 
     vec2 maskTexCoord0 = (vec2(aMaskTexCoord0) + tileOffset) / 256.0;
-    vec2 maskTexCoord1 = (vec2(aMaskTexCoord1) + tileOffset) / 256.0;
 
     vec2 textureMetadataScale = vec2(1.0) / vec2(uTextureMetadataSize);
     vec2 metadataEntryCoord = vec2(aColor % 128 * 4, aColor / 128);
@@ -48,7 +45,6 @@ void main() {
 
     vColorTexCoord0 = mat2(colorTexMatrix0) * position + colorTexOffsets.xy;
     vMaskTexCoord0 = vec3(maskTexCoord0, float(aMaskBackdrop.x));
-    vMaskTexCoord1 = vec3(maskTexCoord1, float(aMaskBackdrop.y));
     vBaseColor = baseColor;
     gl_Position = uTransform * vec4(position, 0.0, 1.0);
 }

shaders/tile_clip.fs.glsl

@@ -0,0 +1,26 @@
+#version 330
+
+// pathfinder/shaders/tile_clip.fs.glsl
+//
+// Copyright © 2020 The Pathfinder Project Developers.
+//
+// Licensed under the Apache License, Version 2.0 <LICENSE-APACHE or
+// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
+// <LICENSE-MIT or http://opensource.org/licenses/MIT>, at your
+// option. This file may not be copied, modified, or distributed
+// except according to those terms.
+
+precision highp float;
+precision highp sampler2D;
+
+uniform sampler2D uSrc;
+
+in vec2 vTexCoord;
+in float vBackdrop;
+
+out vec4 oFragColor;
+
+void main() {
+    float alpha = clamp(abs(texture(uSrc, vTexCoord).r + vBackdrop), 0.0, 1.0);
+    oFragColor = vec4(alpha, 0.0, 0.0, 1.0);
+}

shaders/tile_clip.vs.glsl

@@ -0,0 +1,30 @@
+#version 330
+
+// pathfinder/shaders/tile_clip.vs.glsl
+//
+// Copyright © 2020 The Pathfinder Project Developers.
+//
+// Licensed under the Apache License, Version 2.0 <LICENSE-APACHE or
+// http://www.apache.org/licenses/LICENSE-2.0> or the MIT license
+// <LICENSE-MIT or http://opensource.org/licenses/MIT>, at your
+// option. This file may not be copied, modified, or distributed
+// except according to those terms.
+
+precision highp float;
+precision highp sampler2D;
+
+in ivec2 aTileOffset;
+in ivec2 aDestTileOrigin;
+in ivec2 aSrcTileOrigin;
+in int aSrcBackdrop;
+
+out vec2 vTexCoord;
+out float vBackdrop;
+
+void main() {
+    vec2 destPosition = vec2(aDestTileOrigin + aTileOffset) / vec2(256.0);
+    vec2 srcPosition = vec2(aSrcTileOrigin + aTileOffset) / vec2(256.0);
+    vTexCoord = srcPosition;
+    vBackdrop = float(aSrcBackdrop);
+    gl_Position = vec4(mix(vec2(-1.0), vec2(1.0), destPosition), 0.0, 1.0);
+}