Allow multiple tile pages.

Closes #151.

renderer/src/builder.rs

@@ -12,7 +12,7 @@
 
 use crate::concurrent::executor::Executor;
 use crate::gpu::renderer::{BlendModeExt, MASK_TILES_ACROSS, MASK_TILES_DOWN};
-use crate::gpu_data::{FillBatchPrimitive, RenderCommand, Tile, TileBatch};
+use crate::gpu_data::{AlphaTileId, Fill, FillBatchEntry, RenderCommand, Tile, TileBatch};
 use crate::gpu_data::{TileBatchTexture, TileObjectPrimitive};
 use crate::options::{PreparedBuildOptions, PreparedRenderTransform, RenderCommandListener};
 use crate::paint::{PaintInfo, PaintMetadata};
@@ -33,7 +33,7 @@ use pathfinder_gpu::TextureSamplingFlags;
 use pathfinder_simd::default::{F32x4, I32x4};
 use std::sync::atomic::{AtomicUsize, Ordering};
 use instant::Instant;
-use std::u16;
+use std::u32;
 
 pub(crate) struct SceneBuilder<'a> {
     scene: &'a Scene,
@@ -45,7 +45,7 @@ pub(crate) struct SceneBuilder<'a> {
 #[derive(Debug)]
 pub(crate) struct ObjectBuilder {
     pub built_path: BuiltPath,
-    pub fills: Vec<FillBatchPrimitive>,
+    pub fills: Vec<FillBatchEntry>,
     pub bounds: RectF,
 }
 
@@ -63,17 +63,23 @@ struct BuiltDrawPath {
 #[derive(Debug)]
 pub(crate) struct BuiltPath {
     pub solid_tiles: SolidTiles,
-    pub empty_tiles: Vec<Tile>,
+    pub empty_tiles: Vec<BuiltTile>,
-    pub single_mask_tiles: Vec<Tile>,
+    pub single_mask_tiles: Vec<BuiltTile>,
-    pub dual_mask_tiles: Vec<Tile>,
+    pub dual_mask_tiles: Vec<BuiltTile>,
     pub tiles: DenseTileMap<TileObjectPrimitive>,
     pub fill_rule: FillRule,
 }
 
+#[derive(Clone, Debug)]
+pub struct BuiltTile {
+    pub page: u16,
+    pub tile: Tile,
+}
+
 #[derive(Clone, Debug)]
 pub(crate) enum SolidTiles {
     Occluders(Vec<Occluder>),
-    Regular(Vec<Tile>),
+    Regular(Vec<BuiltTile>),
 }
 
 #[derive(Clone, Copy, Debug)]
@@ -170,8 +176,7 @@ impl<'a> SceneBuilder<'a> {
                                    TilingPathInfo::Clip);
 
         tiler.generate_tiles();
-
+        self.send_fills(tiler.object_builder.fills);
-        self.listener.send(RenderCommand::AddFills(tiler.object_builder.fills));
         tiler.object_builder.built_path
     }
 
@@ -203,9 +208,7 @@ impl<'a> SceneBuilder<'a> {
         }));
 
         tiler.generate_tiles();
-
+        self.send_fills(tiler.object_builder.fills);
-        self.listener.send(RenderCommand::AddFills(tiler.object_builder.fills));
-
         BuiltDrawPath {
             path: tiler.object_builder.built_path,
             blend_mode: path_object.blend_mode(),
@@ -217,6 +220,12 @@ impl<'a> SceneBuilder<'a> {
         }
     }
 
+    fn send_fills(&self, fills: Vec<FillBatchEntry>) {
+        if !fills.is_empty() {
+            self.listener.send(RenderCommand::AddFills(fills));
+        }
+    }
+
     fn cull_tiles(&self, paint_metadata: &[PaintMetadata], built_draw_paths: Vec<BuiltDrawPath>)
                   -> CulledTiles {
         let mut culled_tiles = CulledTiles { display_list: vec![] };
@@ -365,50 +374,98 @@ impl<'a> SceneBuilder<'a> {
     fn add_alpha_tiles(&self,
                        culled_tiles: &mut CulledTiles,
                        layer_z_buffer: &ZBuffer,
-                       alpha_tiles: &[Tile],
+                       built_alpha_tiles: &[BuiltTile],
                        current_depth: u32,
                        color_texture: Option<TileBatchTexture>,
                        blend_mode: BlendMode,
                        filter: Filter,
                        mask_0_fill_rule: Option<FillRule>,
                        mask_1_fill_rule: Option<FillRule>) {
-        if alpha_tiles.is_empty() {
+        let mut batch_indices: Vec<BatchIndex> = vec![];
-            return;
+        for built_alpha_tile in built_alpha_tiles {
-        }
+            // Early cull if possible.
+            let alpha_tile_coords = built_alpha_tile.tile.tile_position();
+            if !layer_z_buffer.test(alpha_tile_coords, current_depth) {
+                continue;
+            }
 
-        // Create a new `DrawTiles` display item if we don't have one or if we have to break a
+            // Find an appropriate batch if we can.
-        // batch due to blend mode or paint page. Note that every path with a blend mode that
+            let mut dest_batch_index = batch_indices.iter().filter(|&batch_index| {
-        // requires a readable framebuffer needs its own batch.
+                batch_index.tile_page == built_alpha_tile.page
-        //
+            }).next().cloned();
-        // TODO(pcwalton): If we really wanted to, we could use tile maps to avoid
+
-        // batch breaks in some cases…
+            // If no batch was found, try to reuse the last batch in the display list.
-        match culled_tiles.display_list.last() {
+            //
-            Some(&CulledDisplayItem::DrawTiles(TileBatch {
+            // TODO(pcwalton): We could try harder to find a batch by taking tile positions into
-                tiles: _,
+            // account...
-                color_texture: ref batch_color_texture,
+            if dest_batch_index.is_none() {
-                blend_mode: batch_blend_mode,
+                match culled_tiles.display_list.last() {
-                filter: batch_filter,
+                    Some(&CulledDisplayItem::DrawTiles(TileBatch {
-                mask_0_fill_rule: batch_mask_0_fill_rule,
+                        tiles: _,
-                mask_1_fill_rule: batch_mask_1_fill_rule,
+                        color_texture: ref batch_color_texture,
-            })) if *batch_color_texture == color_texture &&
+                        blend_mode: batch_blend_mode,
-                batch_blend_mode == blend_mode &&
+                        filter: batch_filter,
-                batch_filter == filter &&
+                        mask_0_fill_rule: batch_mask_0_fill_rule,
-                batch_mask_0_fill_rule == mask_0_fill_rule &&
+                        mask_1_fill_rule: batch_mask_1_fill_rule,
-                batch_mask_1_fill_rule == mask_1_fill_rule &&
+                        tile_page: batch_tile_page
-                !batch_blend_mode.needs_readable_framebuffer() => {}
+                    })) if *batch_color_texture == color_texture &&
-            _ => {
+                            batch_blend_mode == blend_mode &&
-                let batch = TileBatch {
+                            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 {
+                            display_item_index: culled_tiles.display_list.len() - 1,
+                            tile_page: batch_tile_page,
+                        });
+                        batch_indices.push(dest_batch_index.unwrap());
+                    }
+                    _ => {}
+                }
+            }
+
+            // If it's still the case that no suitable batch was found, then make a new one.
+            if dest_batch_index.is_none() {
+                dest_batch_index = Some(BatchIndex {
+                    display_item_index: culled_tiles.display_list.len(),
+                    tile_page: built_alpha_tile.page,
+                });
+                batch_indices.push(dest_batch_index.unwrap());
+                culled_tiles.display_list.push(CulledDisplayItem::DrawTiles(TileBatch {
                     tiles: vec![],
                     color_texture,
                     blend_mode,
                     filter,
                     mask_0_fill_rule,
                     mask_1_fill_rule,
-                };
+                    tile_page: built_alpha_tile.page,
-                culled_tiles.display_list.push(CulledDisplayItem::DrawTiles(batch))
+                }));
+            }
+
+            // Add to the appropriate batch.
+            match culled_tiles.display_list[dest_batch_index.unwrap().display_item_index] {
+                CulledDisplayItem::DrawTiles(ref mut tiles) => {
+                    tiles.tiles.push(built_alpha_tile.tile);
+                }
+                _ => unreachable!(),
             }
         }
 
+        #[derive(Clone, Copy)]
+        struct BatchIndex {
+            display_item_index: usize,
+            tile_page: u16,
+        }
+
+        /*
+        // Create a new `DrawTiles` display item if we don't have one or if we have to break a
+        // batch due to blend mode or paint page. Note that every path with a blend mode that
+        // requires a readable framebuffer needs its own batch.
+        //
+        // TODO(pcwalton): If we really wanted to, we could use tile maps to avoid
+        // batch breaks in some cases…
+
         // Fetch the destination alpha tiles buffer.
         let culled_alpha_tiles = match *culled_tiles.display_list.last_mut().unwrap() {
             CulledDisplayItem::DrawTiles(TileBatch { tiles: ref mut culled_alpha_tiles, .. }) => {
@@ -423,6 +480,7 @@ impl<'a> SceneBuilder<'a> {
                 culled_alpha_tiles.push(*alpha_tile);
             }
         }
+        */
     }
 
     fn pack_tiles(&mut self, culled_tiles: CulledTiles) {
@@ -507,9 +565,9 @@ impl BuiltPath {
         };
 
         BuiltPath {
+            empty_tiles: vec![],
             single_mask_tiles: vec![],
             dual_mask_tiles: vec![],
-            empty_tiles: vec![],
             solid_tiles: if occludes {
                 SolidTiles::Occluders(vec![])
             } else {
@@ -595,8 +653,8 @@ impl ObjectBuilder {
             return;
         }
 
-        // Allocate global tile if necessary.
+        // Allocate a global tile if necessary.
-        let alpha_tile_index = self.get_or_allocate_alpha_tile_index(scene_builder, tile_coords);
+        let alpha_tile_id = self.get_or_allocate_alpha_tile_index(scene_builder, tile_coords);
 
         // Pack whole pixels.
         let px = (segment & I32x4::splat(0xf00)).to_u32x4();
@@ -604,15 +662,18 @@ impl ObjectBuilder {
 
         // Pack instance data.
         debug!("... OK, pushing");
-        self.fills.push(FillBatchPrimitive {
+        self.fills.push(FillBatchEntry {
-            px: LineSegmentU4 { from: px[0] as u8, to: px[2] as u8 },
+            page: alpha_tile_id.page(),
-            subpx: LineSegmentU8 {
+            fill: Fill {
-                from_x: from_x as u8,
+                px: LineSegmentU4 { from: px[0] as u8, to: px[2] as u8 },
-                from_y: from_y as u8,
+                subpx: LineSegmentU8 {
-                to_x:   to_x   as u8,
+                    from_x: from_x as u8,
-                to_y:   to_y   as u8,
+                    from_y: from_y as u8,
+                    to_x:   to_x   as u8,
+                    to_y:   to_y   as u8,
+                },
+                alpha_tile_index: alpha_tile_id.tile(),
             },
-            alpha_tile_index,
         });
     }
 
@@ -620,19 +681,18 @@ impl ObjectBuilder {
         &mut self,
         scene_builder: &SceneBuilder,
         tile_coords: Vector2I,
-    ) -> u16 {
+    ) -> AlphaTileId {
         let local_tile_index = self.built_path.tiles.coords_to_index_unchecked(tile_coords);
-        let alpha_tile_index = self.built_path.tiles.data[local_tile_index].alpha_tile_index;
+        let alpha_tile_id = self.built_path.tiles.data[local_tile_index].alpha_tile_id;
-        if alpha_tile_index != !0 {
+        if alpha_tile_id.is_valid() {
-            return alpha_tile_index;
+            return alpha_tile_id;
         }
 
-        // FIXME(pcwalton): Check for overflow!
+        let alpha_tile_index = scene_builder.next_alpha_tile_index.fetch_add(1, Ordering::Relaxed);
-        let alpha_tile_index = scene_builder
+        debug_assert!(alpha_tile_index < u32::MAX as usize);
-            .next_alpha_tile_index
+        let alpha_tile_id = AlphaTileId(alpha_tile_index as u32);
-            .fetch_add(1, Ordering::Relaxed) as u16;
+        self.built_path.tiles.data[local_tile_index].alpha_tile_id = alpha_tile_id;
-        self.built_path.tiles.data[local_tile_index].alpha_tile_index = alpha_tile_index;
+        alpha_tile_id
-        alpha_tile_index
     }
 
     pub(crate) fn add_active_fill(
@@ -738,21 +798,32 @@ impl ObjectBuilder {
 
 impl<'a> PackedTile<'a> {
     pub(crate) fn add_to(&self,
-                         tiles: &mut Vec<Tile>,
+                         tiles: &mut Vec<BuiltTile>,
                          draw_tiling_path_info: &DrawTilingPathInfo) {
-        let fill_tile_index = self.draw_tile.alpha_tile_index as u16;
+        let fill_tile_index = self.draw_tile.alpha_tile_id.tile() as u16;
         let fill_tile_backdrop = self.draw_tile.backdrop as i8;
         let (clip_tile_index, clip_tile_backdrop) = match self.clip_tile {
             None => (0, 0),
-            Some(clip_tile) => (clip_tile.alpha_tile_index as u16, clip_tile.backdrop as i8),
+            Some(clip_tile) => {
+                // 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)
+            }
         };
 
-        tiles.push(Tile::new_alpha(self.tile_coords,
+        tiles.push(BuiltTile {
-                                   fill_tile_index,
+            page: self.draw_tile.alpha_tile_id.page(),
-                                   fill_tile_backdrop,
+            tile: Tile::new_alpha(self.tile_coords,
-                                   clip_tile_index,
+                                  fill_tile_index,
-                                   clip_tile_backdrop,
+                                  fill_tile_backdrop,
-                                   draw_tiling_path_info));
+                                  clip_tile_index,
+                                  clip_tile_backdrop,
+                                  draw_tiling_path_info),
+        });
     }
 }
 

renderer/src/gpu/renderer.rs

@@ -14,8 +14,9 @@ use crate::gpu::shaders::{BlitProgram, BlitVertexArray, CopyTileProgram, CopyTil
 use crate::gpu::shaders::{FillProgram, FillVertexArray, MAX_FILLS_PER_BATCH, ReprojectionProgram};
 use crate::gpu::shaders::{ReprojectionVertexArray, StencilProgram, StencilVertexArray};
 use crate::gpu::shaders::{TileProgram, TileVertexArray};
-use crate::gpu_data::{FillBatchPrimitive, RenderCommand, TextureLocation, TextureMetadataEntry};
+use crate::gpu_data::{Fill, FillBatchEntry, RenderCommand, TextureLocation};
-use crate::gpu_data::{TexturePageDescriptor, TexturePageId, Tile, TileBatchTexture};
+use crate::gpu_data::{TextureMetadataEntry, TexturePageDescriptor, TexturePageId};
+use crate::gpu_data::{Tile, TileBatchTexture};
 use crate::options::BoundingQuad;
 use crate::paint::PaintCompositeOp;
 use crate::tiles::{TILE_HEIGHT, TILE_WIDTH};
@@ -36,7 +37,6 @@ use pathfinder_gpu::{RenderState, RenderTarget, StencilFunc, StencilState, Textu
 use pathfinder_gpu::{TextureFormat, UniformData};
 use pathfinder_resources::ResourceLoader;
 use pathfinder_simd::default::{F32x2, F32x4, I32x2};
-use std::cmp;
 use std::collections::VecDeque;
 use std::f32;
 use std::mem;
@@ -120,7 +120,7 @@ where
     quads_vertex_indices_buffer: D::Buffer,
     quads_vertex_indices_length: usize,
     fill_vertex_array: FillVertexArray<D>,
-    fill_framebuffer: D::Framebuffer,
+    alpha_tile_pages: Vec<AlphaTilePage<D>>,
     dest_blend_framebuffer: D::Framebuffer,
     intermediate_dest_framebuffer: D::Framebuffer,
     texture_pages: Vec<TexturePage<D>>,
@@ -140,7 +140,6 @@ where
 
     // Rendering state
     framebuffer_flags: FramebufferFlags,
-    buffered_fills: Vec<FillBatchPrimitive>,
     texture_cache: TextureCache<D>,
 
     // Debug
@@ -228,11 +227,6 @@ where
             &quad_vertex_indices_buffer,
         );
 
-        let fill_framebuffer_size = vec2i(MASK_FRAMEBUFFER_WIDTH, MASK_FRAMEBUFFER_HEIGHT);
-        let fill_framebuffer_texture =
-            device.create_texture(TextureFormat::R16F, fill_framebuffer_size);
-        let fill_framebuffer = device.create_framebuffer(fill_framebuffer_texture);
-
         let window_size = dest_framebuffer.window_size(&device);
         let dest_blend_texture = device.create_texture(TextureFormat::RGBA8, window_size);
         let dest_blend_framebuffer = device.create_framebuffer(dest_blend_texture);
@@ -264,7 +258,7 @@ where
             quads_vertex_indices_buffer,
             quads_vertex_indices_length: 0,
             fill_vertex_array,
-            fill_framebuffer,
+            alpha_tile_pages: vec![],
             dest_blend_framebuffer,
             intermediate_dest_framebuffer,
             texture_pages: vec![],
@@ -289,7 +283,6 @@ where
             debug_ui_presenter,
 
             framebuffer_flags: FramebufferFlags::empty(),
-            buffered_fills: vec![],
             texture_cache: TextureCache::new(),
 
             flags: RendererFlags::empty(),
@@ -298,6 +291,10 @@ where
 
     pub fn begin_scene(&mut self) {
         self.framebuffer_flags = FramebufferFlags::empty();
+        for alpha_tile_page in &mut self.alpha_tile_pages {
+            alpha_tile_page.must_preserve_framebuffer = false;
+        }
+
         self.device.begin_commands();
         self.stats = RenderStats::default();
     }
@@ -320,7 +317,11 @@ where
                 self.upload_texture_metadata(metadata)
             }
             RenderCommand::AddFills(ref fills) => self.add_fills(fills),
-            RenderCommand::FlushFills => self.draw_buffered_fills(),
+            RenderCommand::FlushFills => {
+                for page_index in 0..(self.alpha_tile_pages.len() as u16) {
+                    self.draw_buffered_fills(page_index)
+                }
+            }
             RenderCommand::BeginTileDrawing => self.begin_tile_drawing(),
             RenderCommand::PushRenderTarget(render_target_id) => {
                 self.push_render_target(render_target_id)
@@ -330,7 +331,8 @@ where
                 let count = batch.tiles.len();
                 self.stats.alpha_tile_count += count;
                 self.upload_tiles(&batch.tiles);
-                self.draw_tiles(count as u32,
+                self.draw_tiles(batch.tile_page,
+                                count as u32,
                                 batch.color_texture,
                                 batch.mask_0_fill_rule,
                                 batch.mask_1_fill_rule,
@@ -550,44 +552,49 @@ where
         self.quads_vertex_indices_length = length;
     }
 
-    fn add_fills(&mut self, mut fills: &[FillBatchPrimitive]) {
+    fn add_fills(&mut self, fill_batch: &[FillBatchEntry]) {
-        if fills.is_empty() {
+        if fill_batch.is_empty() {
             return;
         }
 
-        self.stats.fill_count += fills.len();
+        self.stats.fill_count += fill_batch.len();
 
-        while !fills.is_empty() {
+        for fill_batch_entry in fill_batch {
-            let count = cmp::min(fills.len(), MAX_FILLS_PER_BATCH - self.buffered_fills.len());
+            let page = fill_batch_entry.page;
-            self.buffered_fills.extend_from_slice(&fills[0..count]);
+            while self.alpha_tile_pages.len() <= page as usize {
-            fills = &fills[count..];
+                self.alpha_tile_pages.push(AlphaTilePage::new(&mut self.device));
-            if self.buffered_fills.len() == MAX_FILLS_PER_BATCH {
+            }
-                self.draw_buffered_fills();
+            self.alpha_tile_pages[page as usize].buffered_fills.push(fill_batch_entry.fill);
+            if self.alpha_tile_pages[page as usize].buffered_fills.len() == MAX_FILLS_PER_BATCH {
+                self.draw_buffered_fills(page);
             }
         }
     }
 
-    fn draw_buffered_fills(&mut self) {
+    fn draw_buffered_fills(&mut self, page: u16) {
-        if self.buffered_fills.is_empty() {
+        let mask_viewport = self.mask_viewport();
+
+        let alpha_tile_page = &mut self.alpha_tile_pages[page as usize];
+        let buffered_fills = &mut alpha_tile_page.buffered_fills;
+        if buffered_fills.is_empty() {
             return;
         }
 
         self.device.allocate_buffer(
             &self.fill_vertex_array.vertex_buffer,
-            BufferData::Memory(&self.buffered_fills),
+            BufferData::Memory(&buffered_fills),
             BufferTarget::Vertex,
             BufferUploadMode::Dynamic,
         );
 
         let mut clear_color = None;
-        if !self.framebuffer_flags.contains(
+        if !alpha_tile_page.must_preserve_framebuffer {
-                FramebufferFlags::MUST_PRESERVE_FILL_FRAMEBUFFER_CONTENTS) {
             clear_color = Some(ColorF::default());
         };
 
-        debug_assert!(self.buffered_fills.len() <= u32::MAX as usize);
+        debug_assert!(buffered_fills.len() <= u32::MAX as usize);
-        self.device.draw_elements_instanced(6, self.buffered_fills.len() as u32, &RenderState {
+        self.device.draw_elements_instanced(6, buffered_fills.len() as u32, &RenderState {
-            target: &RenderTarget::Framebuffer(&self.fill_framebuffer),
+            target: &RenderTarget::Framebuffer(&alpha_tile_page.framebuffer),
             program: &self.fill_program.program,
             vertex_array: &self.fill_vertex_array.vertex_array,
             primitive: Primitive::Triangles,
@@ -600,7 +607,7 @@ where
                  UniformData::Vec2(F32x2::new(TILE_WIDTH as f32, TILE_HEIGHT as f32))),
                 (&self.fill_program.area_lut_uniform, UniformData::TextureUnit(0)),
             ],
-            viewport: self.mask_viewport(),
+            viewport: mask_viewport,
             options: RenderOptions {
                 blend: Some(BlendState {
                     src_rgb_factor: BlendFactor::One,
@@ -614,8 +621,8 @@ where
             },
         });
 
-        self.framebuffer_flags.insert(FramebufferFlags::MUST_PRESERVE_FILL_FRAMEBUFFER_CONTENTS);
+        alpha_tile_page.must_preserve_framebuffer = true;
-        self.buffered_fills.clear();
+        buffered_fills.clear();
     }
 
     fn tile_transform(&self) -> Transform4F {
@@ -625,6 +632,7 @@ where
     }
 
     fn draw_tiles(&mut self,
+                  tile_page: u16,
                   tile_count: u32,
                   color_texture_0: Option<TileBatchTexture>,
                   mask_0_fill_rule: Option<FillRule>,
@@ -676,12 +684,14 @@ where
         if mask_0_fill_rule.is_some() {
             uniforms.push((&self.tile_program.mask_texture_0_uniform,
                            UniformData::TextureUnit(textures.len() as u32)));
-            textures.push(self.device.framebuffer_texture(&self.fill_framebuffer));
+            textures.push(self.device.framebuffer_texture(
+                &self.alpha_tile_pages[tile_page as usize].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.fill_framebuffer));
+            textures.push(self.device.framebuffer_texture(
+                &self.alpha_tile_pages[tile_page as usize].framebuffer));
         }
 
         // TODO(pcwalton): Refactor.
@@ -1158,9 +1168,8 @@ impl Div<usize> for RenderTime {
 
 bitflags! {
     struct FramebufferFlags: u8 {
-        const MUST_PRESERVE_FILL_FRAMEBUFFER_CONTENTS = 0x01;
+        const MUST_PRESERVE_MASK_FRAMEBUFFER_CONTENTS = 0x01;
-        const MUST_PRESERVE_MASK_FRAMEBUFFER_CONTENTS = 0x02;
+        const MUST_PRESERVE_DEST_FRAMEBUFFER_CONTENTS = 0x02;
-        const MUST_PRESERVE_DEST_FRAMEBUFFER_CONTENTS = 0x04;
     }
 }
 
@@ -1369,6 +1378,21 @@ impl BlendModeExt for BlendMode {
     }
 }
 
+struct AlphaTilePage<D> where D: Device {
+    buffered_fills: Vec<Fill>,
+    framebuffer: D::Framebuffer,
+    must_preserve_framebuffer: bool,
+}
+
+impl<D> AlphaTilePage<D> where D: Device {
+    fn new(device: &mut D) -> AlphaTilePage<D> {
+        let framebuffer_size = vec2i(MASK_FRAMEBUFFER_WIDTH, MASK_FRAMEBUFFER_HEIGHT);
+        let framebuffer_texture = device.create_texture(TextureFormat::R16F, framebuffer_size);
+        let framebuffer = device.create_framebuffer(framebuffer_texture);
+        AlphaTilePage { buffered_fills: vec![], framebuffer, must_preserve_framebuffer: false }
+    }
+}
+
 bitflags! {
     struct RendererFlags: u8 {
         // Whether we need a depth buffer.

renderer/src/gpu/shaders.rs

@@ -8,7 +8,7 @@
 // option. This file may not be copied, modified, or distributed
 // except according to those terms.
 
-use crate::gpu_data::FillBatchPrimitive;
+use crate::gpu_data::Fill;
 use pathfinder_gpu::{BufferData, BufferTarget, BufferUploadMode, Device, VertexAttrClass};
 use pathfinder_gpu::{VertexAttrDescriptor, VertexAttrType};
 use pathfinder_resources::ResourceLoader;
@@ -69,8 +69,7 @@ where
         let vertex_array = device.create_vertex_array();
 
         let vertex_buffer = device.create_buffer();
-        let vertex_buffer_data: BufferData<FillBatchPrimitive> =
+        let vertex_buffer_data: BufferData<Fill> = BufferData::Uninitialized(MAX_FILLS_PER_BATCH);
-            BufferData::Uninitialized(MAX_FILLS_PER_BATCH);
         device.allocate_buffer(
             &vertex_buffer,
             vertex_buffer_data,

renderer/src/gpu_data.rs

@@ -56,7 +56,7 @@ pub enum RenderCommand {
     UploadTextureMetadata(Vec<TextureMetadataEntry>),
 
     // Adds fills to the queue.
-    AddFills(Vec<FillBatchPrimitive>),
+    AddFills(Vec<FillBatchEntry>),
 
     // Flushes the queue of fills.
     FlushFills,
@@ -100,6 +100,7 @@ pub struct TileBatch {
     pub mask_1_fill_rule: Option<FillRule>,
     pub filter: Filter,
     pub blend_mode: BlendMode,
+    pub tile_page: u16,
 }
 
 #[derive(Clone, Copy, Debug, PartialEq)]
@@ -120,8 +121,7 @@ pub struct FillObjectPrimitive {
 #[derive(Clone, Copy, Debug)]
 #[repr(C)]
 pub struct TileObjectPrimitive {
-    /// If `u16::MAX`, then this is a solid tile.
+    pub alpha_tile_id: AlphaTileId,
-    pub alpha_tile_index: u16,
     pub backdrop: i8,
 }
 
@@ -132,10 +132,16 @@ pub struct TextureMetadataEntry {
     pub base_color: ColorU,
 }
 
+#[derive(Clone, Copy, Debug, Default)]
+pub struct FillBatchEntry {
+    pub fill: Fill,
+    pub page: u16,
+}
+
 // FIXME(pcwalton): Move `subpx` before `px` and remove `repr(packed)`.
 #[derive(Clone, Copy, Debug, Default)]
 #[repr(packed)]
-pub struct FillBatchPrimitive {
+pub struct Fill {
     pub px: LineSegmentU4,
     pub subpx: LineSegmentU8,
     pub alpha_tile_index: u16,
@@ -155,6 +161,31 @@ pub struct Tile {
     pub color: u16,
 }
 
+#[derive(Clone, Copy, PartialEq, Debug)]
+pub struct AlphaTileId(pub u32);
+
+impl AlphaTileId {
+    #[inline]
+    pub fn invalid() -> AlphaTileId {
+        AlphaTileId(!0)
+    }
+
+    #[inline]
+    pub fn page(self) -> u16 {
+        (self.0 >> 16) as u16
+    }
+
+    #[inline]
+    pub fn tile(self) -> u16 {
+        (self.0 & 0xffff) as u16
+    }
+
+    #[inline]
+    pub fn is_valid(self) -> bool {
+        self.0 < !0
+    }
+}
+
 impl Debug for RenderCommand {
     fn fmt(&self, formatter: &mut Formatter) -> DebugResult {
         match *self {
@@ -171,7 +202,9 @@ impl Debug for RenderCommand {
             RenderCommand::UploadTextureMetadata(ref metadata) => {
                 write!(formatter, "UploadTextureMetadata(x{})", metadata.len())
             }
-            RenderCommand::AddFills(ref fills) => write!(formatter, "AddFills(x{})", fills.len()),
+            RenderCommand::AddFills(ref fills) => {
+                write!(formatter, "AddFills(x{})", fills.len())
+            }
             RenderCommand::FlushFills => write!(formatter, "FlushFills"),
             RenderCommand::PushRenderTarget(render_target_id) => {
                 write!(formatter, "PushRenderTarget({:?})", render_target_id)

renderer/src/tiles.rs

@@ -9,7 +9,7 @@
 // except according to those terms.
 
 use crate::builder::{BuiltPath, ObjectBuilder, Occluder, SceneBuilder, SolidTiles};
-use crate::gpu_data::TileObjectPrimitive;
+use crate::gpu_data::{AlphaTileId, TileObjectPrimitive};
 use crate::paint::{PaintId, PaintMetadata};
 use pathfinder_content::effects::BlendMode;
 use pathfinder_content::fill::FillRule;
@@ -151,10 +151,12 @@ impl<'a> Tiler<'a> {
                     }
                 }
                 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);
                 }
                 TileType::DualMask => {
+                    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);
                 }
@@ -413,6 +415,30 @@ impl<'a> PackedTile<'a> {
            -> PackedTile<'a> {
         let tile_coords = object_builder.local_tile_index_to_coords(draw_tile_index as u32);
 
+        // First, if the draw tile is empty, cull it regardless of clip.
+        if draw_tile.is_solid() {
+            match (object_builder.built_path.fill_rule, draw_tile.backdrop) {
+                (FillRule::Winding, 0) => {
+                    return PackedTile {
+                        tile_type: TileType::Empty,
+                        tile_coords,
+                        draw_tile,
+                        clip_tile: None,
+                    };
+                }
+                (FillRule::Winding, _) => {}
+                (FillRule::EvenOdd, backdrop) if backdrop % 2 == 0 => {
+                    return PackedTile {
+                        tile_type: TileType::Empty,
+                        tile_coords,
+                        draw_tile,
+                        clip_tile: None,
+                    };
+                }
+                (FillRule::EvenOdd, _) => {}
+            }
+        }
+
         // Figure out what clip tile we need, if any.
         let clip_tile = match draw_tiling_path_info.built_clip_path {
             None => None,
@@ -448,51 +474,43 @@ impl<'a> PackedTile<'a> {
             }
         };
 
-        if clip_tile.is_none() {
+        // Choose a tile type.
-            if draw_tile.is_solid() {
+        match clip_tile {
-                // This is the simple case of a solid tile with no clip, so there are optimization
+            None if draw_tile.is_solid() => {
-                // opportunities. First, tiles that must be blank per the fill rule are always
+                // This is a solid tile that completely occludes the background.
-                // skipped.
+                PackedTile { tile_type: TileType::Solid, tile_coords, draw_tile, clip_tile }
-                match (object_builder.built_path.fill_rule, draw_tile.backdrop) {
+            }
-                    (FillRule::Winding, 0) => {
+            None => {
-                        return PackedTile {
+                // We have a draw tile and no clip tile.
-                            tile_type: TileType::Empty,
+                PackedTile {
-                            tile_coords,
+                    tile_type: TileType::SingleMask,
-                            draw_tile,
-                            clip_tile,
-                        };
-                    }
-                    (FillRule::Winding, _) => {}
-                    (FillRule::EvenOdd, backdrop) if backdrop % 2 == 0 => {
-                        return PackedTile {
-                            tile_type: TileType::Empty,
-                            tile_coords,
-                            draw_tile,
-                            clip_tile,
-                        };
-                    }
-                    (FillRule::EvenOdd, _) => {}
-                }
-
-                // Next, if this is a solid tile that completely occludes the background, record
-                // that fact. Otherwise, add a regular solid tile.
-                return PackedTile {
-                    tile_type: TileType::Solid,
                     tile_coords,
                     draw_tile,
-                    clip_tile,
+                    clip_tile: None,
-                };
+                }
+            }
+            Some(clip_tile) if draw_tile.is_solid() => {
+                // We have a solid draw tile and a clip tile. This is effectively the same as
+                // having a draw tile and no clip tile.
+                //
+                // FIXME(pcwalton): This doesn't preserve the fill rule of the clip path!
+                PackedTile {
+                    tile_type: TileType::SingleMask,
+                    tile_coords,
+                    draw_tile: clip_tile,
+                    clip_tile: None,
+                }
+            }
+            Some(clip_tile) => {
+                // We have both a draw and clip mask. Composite them together.
+                PackedTile {
+                    tile_type: TileType::DualMask,
+                    tile_coords,
+                    draw_tile,
+                    clip_tile: Some(clip_tile),
+                }
             }
-
-            return PackedTile {
-                tile_type: TileType::SingleMask,
-                tile_coords,
-                draw_tile,
-                clip_tile,
-            };
         }
-
-        PackedTile { tile_type: TileType::DualMask, tile_coords, draw_tile, clip_tile }
     }
 }
 
@@ -671,11 +689,11 @@ impl PartialOrd<ActiveEdge> for ActiveEdge {
 impl Default for TileObjectPrimitive {
     #[inline]
     fn default() -> TileObjectPrimitive {
-        TileObjectPrimitive { backdrop: 0, alpha_tile_index: !0 }
+        TileObjectPrimitive { backdrop: 0, alpha_tile_id: AlphaTileId::invalid() }
     }
 }
 
 impl TileObjectPrimitive {
     #[inline]
-    pub fn is_solid(&self) -> bool { self.alpha_tile_index == !0 }
+    pub fn is_solid(&self) -> bool { !self.alpha_tile_id.is_valid() }
 }

renderer/src/z_buffer.rs

@@ -93,6 +93,7 @@ impl ZBuffer {
                         blend_mode: BlendMode::default(),
                         mask_0_fill_rule: None,
                         mask_1_fill_rule: None,
+                        tile_page: !0,
                     });
                 }
             }