diff --git a/renderer/src/builder.rs b/renderer/src/builder.rs index 07ae83fe..5f1ac45c 100644 --- a/renderer/src/builder.rs +++ b/renderer/src/builder.rs @@ -14,11 +14,10 @@ use crate::concurrent::executor::Executor; use crate::gpu::blend::BlendModeExt; use crate::gpu::options::RendererLevel; use crate::gpu_data::{AlphaTileId, BackdropInfoD3D11, Clip, ClippedPathInfo, DiceMetadataD3D11}; -use crate::gpu_data::{DrawTileBatch, DrawTileBatchD3D9, DrawTileBatchD3D11, Fill, PathBatchIndex}; -use crate::gpu_data::{PathSource, PrepareTilesInfoD3D11}; -use crate::gpu_data::{PropagateMetadataD3D11, RenderCommand, SegmentIndicesD3D11, SegmentsD3D11}; -use crate::gpu_data::{TileBatchDataD3D11, TileBatchId, TileBatchTexture}; -use crate::gpu_data::{TileObjectPrimitive, TilePathInfoD3D11}; +use crate::gpu_data::{DrawTileBatch, DrawTileBatchD3D9, DrawTileBatchD3D11, Fill, GlobalPathId}; +use crate::gpu_data::{PathBatchIndex, PathSource, PrepareTilesInfoD3D11, PropagateMetadataD3D11}; +use crate::gpu_data::{RenderCommand, SegmentIndicesD3D11, SegmentsD3D11, TileBatchDataD3D11}; +use crate::gpu_data::{TileBatchId, TileBatchTexture, TileObjectPrimitive, TilePathInfoD3D11}; use crate::options::{PrepareMode, PreparedBuildOptions, PreparedRenderTransform}; use crate::paint::{PaintId, PaintInfo, PaintMetadata}; use crate::scene::{ClipPathId, DisplayItem, DrawPath, DrawPathId, LastSceneInfo, PathId}; @@ -48,6 +47,8 @@ pub(crate) const ALPHA_TILES_PER_LEVEL: usize = 1 << (32 - ALPHA_TILE_LEVEL_COUN const CURVE_IS_QUADRATIC: u32 = 0x80000000; const CURVE_IS_CUBIC: u32 = 0x40000000; +const MAX_CLIP_BATCHES: u32 = 32; + pub(crate) struct SceneBuilder<'a, 'b, 'c, 'd> { scene: &'a mut Scene, built_options: &'b PreparedBuildOptions, @@ -265,6 +266,7 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> { path_object.fill_rule(), view_box, &prepare_mode, + path_object.clip_path(), &[], TilingPathInfo::Clip); @@ -298,11 +300,11 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> { path_object.fill_rule(), view_box, &prepare_mode, + path_object.clip_path(), &built_clip_paths, TilingPathInfo::Draw(DrawTilingPathInfo { paint_id, blend_mode: path_object.blend_mode(), - clip_path_id: path_object.clip_path(), fill_rule: path_object.fill_rule(), })); @@ -322,7 +324,7 @@ impl<'a, 'b, 'c, 'd> SceneBuilder<'a, 'b, 'c, 'd> { paint_metadata: &[PaintMetadata], prepare_mode: &PrepareMode, built_paths: Option) { - let mut tile_batch_builder = TileBatchBuilder::new(&prepare_mode, built_paths); + let mut tile_batch_builder = TileBatchBuilder::new(built_paths); // Prepare display items. for display_item in self.scene.display_list() { @@ -410,6 +412,7 @@ impl BuiltPath { view_box_bounds: RectF, fill_rule: FillRule, prepare_mode: &PrepareMode, + clip_path_id: Option, tiling_path_info: &TilingPathInfo) -> BuiltPath { let paint_id = match *tiling_path_info { @@ -427,13 +430,6 @@ impl BuiltPath { let tile_bounds = tiles::round_rect_out_to_tile_bounds(tile_map_bounds); - let clip_path_id = match *tiling_path_info { - TilingPathInfo::Draw(ref draw_tiling_path_info) => { - draw_tiling_path_info.clip_path_id - } - _ => None, - }; - let data = match *prepare_mode { PrepareMode::CPU => { BuiltPathData::CPU(BuiltPathBinCPUData { @@ -492,6 +488,7 @@ impl ObjectBuilder { view_box_bounds: RectF, fill_rule: FillRule, prepare_mode: &PrepareMode, + clip_path_id: Option, tiling_path_info: &TilingPathInfo) -> ObjectBuilder { let built_path = BuiltPath::new(path_id, @@ -499,6 +496,7 @@ impl ObjectBuilder { view_box_bounds, fill_rule, prepare_mode, + clip_path_id, tiling_path_info); ObjectBuilder { built_path, bounds: path_bounds, fills: vec![] } } @@ -650,7 +648,7 @@ impl TileBatchDataD3D11 { fn push(&mut self, path: &BuiltPath, global_path_id: PathId, - batch_clip_path_index: Option, + batch_clip_path_id: Option, z_write: bool, sink: &SceneSink) -> PathBatchIndex { @@ -662,7 +660,10 @@ impl TileBatchDataD3D11 { tile_offset: self.tile_count, path_index: batch_path_index, z_write: z_write as u32, - clip_path_index: batch_clip_path_index.unwrap_or(PathBatchIndex::none()), + clip_path_index: match batch_clip_path_id { + None => PathBatchIndex::none(), + Some(batch_clip_path_id) => batch_clip_path_id.path_index, + }, backdrop_offset: self.prepare_info.backdrops.len() as u32, pad0: 0, pad1: 0, @@ -716,13 +717,14 @@ impl TileBatchDataD3D11 { // Handle clip. - if batch_clip_path_index.is_none() { - return batch_path_index; - } + let clip_batch_id = match batch_clip_path_id { + None => return batch_path_index, + Some(batch_clip_path_id) => batch_clip_path_id.batch_id, + }; if self.clipped_path_info.is_none() { self.clipped_path_info = Some(ClippedPathInfo { - clip_batch_id: TileBatchId(0), + clip_batch_id, clipped_path_count: 0, max_clipped_tile_count: 0, clips: match sink.renderer_level { @@ -836,31 +838,33 @@ impl SegmentsD3D11 { struct TileBatchBuilder { prepare_commands: Vec, draw_commands: Vec, - clip_id_to_path_batch_index: FxHashMap, + clip_batches_d3d11: Option, next_batch_id: TileBatchId, level: TileBatchBuilderLevel, } enum TileBatchBuilderLevel { D3D9 { built_paths: BuiltPaths }, - D3D11 { clip_prepare_batch: TileBatchDataD3D11 }, + D3D11, } impl TileBatchBuilder { - fn new(prepare_mode: &PrepareMode, built_paths: Option) -> TileBatchBuilder { + fn new(built_paths: Option) -> TileBatchBuilder { TileBatchBuilder { prepare_commands: vec![], draw_commands: vec![], - next_batch_id: TileBatchId(1), - clip_id_to_path_batch_index: FxHashMap::default(), - level: match built_paths { + next_batch_id: TileBatchId(MAX_CLIP_BATCHES), + clip_batches_d3d11: match built_paths { None => { - TileBatchBuilderLevel::D3D11 { - clip_prepare_batch: TileBatchDataD3D11::new(TileBatchId(0), - &prepare_mode, - PathSource::Clip), - } + Some(ClipBatchesD3D11 { + prepare_batches: vec![], + clip_id_to_path_batch_index: FxHashMap::default(), + }) } + Some(_) => None, + }, + level: match built_paths { + None => TileBatchBuilderLevel::D3D11, Some(built_paths) => TileBatchBuilderLevel::D3D9 { built_paths }, }, } @@ -945,34 +949,16 @@ impl TileBatchBuilder { } // Add clip path if necessary. - let clip_path = match draw_path.clip_path_id { + let clip_path = match self.clip_batches_d3d11 { None => None, - Some(clip_path_id) => { - match self.clip_id_to_path_batch_index.get(&clip_path_id) { - Some(&clip_path_batch_index) => Some(clip_path_batch_index), - None => { - match self.level { - TileBatchBuilderLevel::D3D9 { .. } => None, - TileBatchBuilderLevel::D3D11 { ref mut clip_prepare_batch } => { - let clip_path = - prepare_clip_path_for_gpu_binning(scene, - built_options, - clip_path_id, - prepare_mode); - let clip_path_index = - clip_prepare_batch.push(&clip_path, - clip_path_id.to_path_id(), - None, - true, - sink); - self.clip_id_to_path_batch_index.insert(clip_path_id, - clip_path_index); - Some(clip_path_index) - } - } - - } - } + Some(ref mut clip_batches_d3d11) => { + add_clip_path_to_batch(scene, + sink, + built_options, + draw_path.clip_path_id, + prepare_mode, + 0, + clip_batches_d3d11) } }; @@ -1071,19 +1057,21 @@ impl TileBatchBuilder { effective_view_box, draw_path.fill_rule(), &prepare_mode, + draw_path.clip_path(), &TilingPathInfo::Draw(DrawTilingPathInfo { paint_id, blend_mode: draw_path.blend_mode(), - clip_path_id: draw_path.clip_path(), fill_rule: draw_path.fill_rule(), })); Some(BuiltDrawPath::new(built_path, draw_path, paint_metadata)) } fn send_to(self, sink: &SceneSink) { - if let TileBatchBuilderLevel::D3D11 { clip_prepare_batch } = self.level { - if clip_prepare_batch.path_count > 0 { - sink.listener.send(RenderCommand::PrepareClipTilesD3D11(clip_prepare_batch)); + if let Some(clip_batches_d3d11) = self.clip_batches_d3d11 { + for prepare_batch in clip_batches_d3d11.prepare_batches.into_iter().rev() { + if prepare_batch.path_count > 0 { + sink.listener.send(RenderCommand::PrepareClipTilesD3D11(prepare_batch)); + } } } @@ -1096,11 +1084,76 @@ impl TileBatchBuilder { } } +struct ClipBatchesD3D11 { + // Will be submitted in reverse (LIFO) order. + prepare_batches: Vec, + clip_id_to_path_batch_index: FxHashMap, +} + +fn add_clip_path_to_batch(scene: &Scene, + sink: &SceneSink, + built_options: &PreparedBuildOptions, + clip_path_id: Option, + prepare_mode: &PrepareMode, + clip_level: usize, + clip_batches_d3d11: &mut ClipBatchesD3D11) + -> Option { + match clip_path_id { + None => None, + Some(clip_path_id) => { + match clip_batches_d3d11.clip_id_to_path_batch_index.get(&clip_path_id) { + Some(&clip_path_batch_index) => { + Some(GlobalPathId { + batch_id: TileBatchId(clip_level as u32), + path_index: clip_path_batch_index, + }) + } + None => { + let PreparedClipPath { + built_path: clip_path, + subclip_id + } = prepare_clip_path_for_gpu_binning(scene, + sink, + built_options, + clip_path_id, + prepare_mode, + clip_level, + clip_batches_d3d11); + while clip_level >= clip_batches_d3d11.prepare_batches.len() { + let clip_tile_batch_id = + TileBatchId(clip_batches_d3d11.prepare_batches.len() as u32); + clip_batches_d3d11.prepare_batches + .push(TileBatchDataD3D11::new(clip_tile_batch_id, + &prepare_mode, + PathSource::Clip)); + } + let clip_path_batch_index = + clip_batches_d3d11.prepare_batches[clip_level] + .push(&clip_path, + clip_path_id.to_path_id(), + subclip_id, + true, + sink); + clip_batches_d3d11.clip_id_to_path_batch_index.insert(clip_path_id, + clip_path_batch_index); + Some(GlobalPathId { + batch_id: TileBatchId(clip_level as u32), + path_index: clip_path_batch_index, + }) + } + } + } + } +} + fn prepare_clip_path_for_gpu_binning(scene: &Scene, + sink: &SceneSink, built_options: &PreparedBuildOptions, clip_path_id: ClipPathId, - prepare_mode: &PrepareMode) - -> BuiltPath { + prepare_mode: &PrepareMode, + clip_level: usize, + clip_batches_d3d11: &mut ClipBatchesD3D11) + -> PreparedClipPath { let transform = match *prepare_mode { PrepareMode::GPU { transform } => transform, PrepareMode::CPU | PrepareMode::TransformCPUBinGPU => { @@ -1109,14 +1162,34 @@ fn prepare_clip_path_for_gpu_binning(scene: &Scene, }; let effective_view_box = scene.effective_view_box(built_options); let clip_path = scene.get_clip_path(clip_path_id); + + // Add subclip path if necessary. + let subclip_id = add_clip_path_to_batch(scene, + sink, + built_options, + clip_path.clip_path(), + prepare_mode, + clip_level + 1, + clip_batches_d3d11); + let path_bounds = transform * clip_path.outline().bounds(); + // TODO(pcwalton): Clip to view box! - BuiltPath::new(clip_path_id.to_path_id(), - path_bounds, - effective_view_box, - clip_path.fill_rule(), - &prepare_mode, - &TilingPathInfo::Clip) + + let built_path = BuiltPath::new(clip_path_id.to_path_id(), + path_bounds, + effective_view_box, + clip_path.fill_rule(), + &prepare_mode, + clip_path.clip_path(), + &TilingPathInfo::Clip); + + PreparedClipPath { built_path, subclip_id } +} + +struct PreparedClipPath { + built_path: BuiltPath, + subclip_id: Option, } fn fixup_batch_for_new_path_if_possible(batch_color_texture: &mut Option, diff --git a/renderer/src/gpu_data.rs b/renderer/src/gpu_data.rs index 9c14fd3f..053f6067 100644 --- a/renderer/src/gpu_data.rs +++ b/renderer/src/gpu_data.rs @@ -179,8 +179,6 @@ pub struct SegmentIndicesD3D11 { #[derive(Clone, Debug)] pub struct ClippedPathInfo { /// The ID of the batch containing the clips. - /// - /// In the current implementation, this is always 0. pub clip_batch_id: TileBatchId, /// The number of paths that have clips. @@ -205,6 +203,12 @@ pub struct PathBatchIndex(pub u32); #[derive(Clone, Copy, PartialEq, Debug)] pub struct TileBatchId(pub u32); +#[derive(Clone, Copy, PartialEq, Debug)] +pub struct GlobalPathId { + pub batch_id: TileBatchId, + pub path_index: PathBatchIndex, +} + #[derive(Clone, Debug)] pub enum DrawTileBatch { D3D9(DrawTileBatchD3D9), @@ -433,6 +437,13 @@ impl PathBatchIndex { } } +impl GlobalPathId { + #[inline] + pub fn none() -> GlobalPathId { + GlobalPathId { batch_id: TileBatchId(!0), path_index: PathBatchIndex(!0) } + } +} + impl AlphaTileId { #[inline] pub fn new(next_alpha_tile_index: &[AtomicUsize; ALPHA_TILE_LEVEL_COUNT], level: usize) diff --git a/renderer/src/scene.rs b/renderer/src/scene.rs index 0d05e131..d536fe01 100644 --- a/renderer/src/scene.rs +++ b/renderer/src/scene.rs @@ -352,6 +352,7 @@ pub struct DrawPath { #[derive(Clone, Debug)] pub struct ClipPath { pub outline: Outline, + pub clip_path: Option, pub fill_rule: FillRule, pub name: String, } @@ -447,7 +448,7 @@ impl DrawPath { impl ClipPath { #[inline] pub fn new(outline: Outline) -> ClipPath { - ClipPath { outline, fill_rule: FillRule::Winding, name: String::new() } + ClipPath { outline, clip_path: None, fill_rule: FillRule::Winding, name: String::new() } } #[inline] @@ -455,6 +456,16 @@ impl ClipPath { &self.outline } + #[inline] + pub(crate) fn clip_path(&self) -> Option { + self.clip_path + } + + #[inline] + pub fn set_clip_path(&mut self, new_clip_path: Option) { + self.clip_path = new_clip_path + } + #[inline] pub(crate) fn fill_rule(&self) -> FillRule { self.fill_rule diff --git a/renderer/src/tiler.rs b/renderer/src/tiler.rs index cb0dcf6a..f2a088a2 100644 --- a/renderer/src/tiler.rs +++ b/renderer/src/tiler.rs @@ -15,8 +15,8 @@ use crate::builder::{BuiltPath, BuiltPathBinCPUData, BuiltPathData, ObjectBuilde use crate::gpu::options::RendererLevel; use crate::gpu_data::AlphaTileId; use crate::options::PrepareMode; -use crate::scene::PathId; -use crate::tiles::{DrawTilingPathInfo, TILE_HEIGHT, TILE_WIDTH, TilingPathInfo}; +use crate::scene::{ClipPathId, PathId}; +use crate::tiles::{TILE_HEIGHT, TILE_WIDTH, TilingPathInfo}; use pathfinder_content::fill::FillRule; use pathfinder_content::outline::{ContourIterFlags, Outline}; use pathfinder_content::segment::Segment; @@ -41,15 +41,14 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> { fill_rule: FillRule, view_box: RectF, prepare_mode: &PrepareMode, + clip_path_id: Option, built_clip_paths: &'a [BuiltPath], path_info: TilingPathInfo) -> Tiler<'a, 'b, 'c, 'd> { let bounds = outline.bounds().intersection(view_box).unwrap_or(RectF::default()); - let clip_path = match path_info { - TilingPathInfo::Draw(DrawTilingPathInfo { clip_path_id: Some(clip_path_id), .. }) => { - Some(&built_clip_paths[clip_path_id.0 as usize]) - } + let clip_path = match clip_path_id { + Some(clip_path_id) => Some(&built_clip_paths[clip_path_id.0 as usize]), _ => None, }; @@ -58,6 +57,7 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> { view_box, fill_rule, prepare_mode, + clip_path_id, &path_info); Tiler { scene_builder, object_builder, outline, clip_path } @@ -158,41 +158,6 @@ impl<'a, 'b, 'c, 'd> Tiler<'a, 'b, 'c, 'd> { backdrops[column] += delta; } - - /* - - // Calculate clips. - let built_clip_path = match self.path_info { - TilingPathInfo::Draw(DrawTilingPathInfo { - built_clip_path: Some(built_clip_path), - .. - }) => built_clip_path, - _ => return, - }; - - let clip_tiles = self.object_builder - .built_path - .clip_tiles - .as_mut() - .expect("Where are the clip tiles?"); - - for draw_tile in &mut self.object_builder.built_path.tiles.data { - let tile_coords = vec2i(draw_tile.tile_x as i32, draw_tile.tile_y as i32); - let built_clip_tile = match built_clip_path.tiles.get(tile_coords) { - None => { - draw_tile.alpha_tile_id = AlphaTileId(!0); - continue; - } - Some(built_clip_tile) => built_clip_tile, - }; - - let clip_tile = clip_tiles.get_mut(tile_coords).unwrap(); - clip_tile.dest_tile_id = draw_tile.alpha_tile_id; - clip_tile.dest_backdrop = draw_tile.backdrop as i32; - clip_tile.src_tile_id = built_clip_tile.alpha_tile_id; - clip_tile.src_backdrop = built_clip_tile.backdrop as i32; - } - */ } } diff --git a/renderer/src/tiles.rs b/renderer/src/tiles.rs index a56406c2..e1a1bc61 100644 --- a/renderer/src/tiles.rs +++ b/renderer/src/tiles.rs @@ -31,7 +31,6 @@ pub(crate) struct DrawTilingPathInfo { pub(crate) paint_id: PaintId, pub(crate) blend_mode: BlendMode, pub(crate) fill_rule: FillRule, - pub(crate) clip_path_id: Option, } impl TilingPathInfo { diff --git a/resources/shaders/gl4/d3d11/propagate.cs.glsl b/resources/shaders/gl4/d3d11/propagate.cs.glsl index a4d9d182..538e7532 100644 --- a/resources/shaders/gl4/d3d11/propagate.cs.glsl +++ b/resources/shaders/gl4/d3d11/propagate.cs.glsl @@ -150,11 +150,6 @@ void main(){ clipTileRect, clipTileCoord); - - - - - int thisClipAlphaTileIndex = int(iClipTiles[clipTileIndex * 4 + 2]<< 8)>> 8; @@ -186,8 +181,6 @@ void main(){ drawTileBackdrop = 0; needNewAlphaTile = false; - } else { - needNewAlphaTile = true; } } } else { diff --git a/resources/shaders/metal/d3d11/propagate.cs.metal b/resources/shaders/metal/d3d11/propagate.cs.metal index 9cb008c6..60af665d 100644 --- a/resources/shaders/metal/d3d11/propagate.cs.metal +++ b/resources/shaders/metal/d3d11/propagate.cs.metal @@ -61,7 +61,7 @@ uint calculateTileIndex(thread const uint& bufferOffset, thread const uint4& til return (bufferOffset + (tileCoord.y * (tileRect.z - tileRect.x))) + tileCoord.x; } -kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bIndirectDrawParams& _303 [[buffer(6)]], device bAlphaTiles& _310 [[buffer(7)]], device bZBuffer& _381 [[buffer(10)]], device bFirstTileMap& _398 [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirstAlphaTileIndex [[buffer(8)]], constant int2& uFramebufferTileSize [[buffer(9)]], const device bBackdrops& _59 [[buffer(1)]], const device bDrawMetadata& _85 [[buffer(2)]], const device bClipMetadata& _126 [[buffer(3)]], device bDrawTiles& _175 [[buffer(4)]], device bClipTiles& _252 [[buffer(5)]], device bIndirectDrawParams& _302 [[buffer(6)]], device bAlphaTiles& _309 [[buffer(7)]], device bZBuffer& _380 [[buffer(10)]], device bFirstTileMap& _397 [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint columnIndex = gl_GlobalInvocationID.x; if (int(columnIndex) >= uColumnCount) @@ -144,10 +144,6 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst drawTileBackdrop = 0; needNewAlphaTile = false; } - else - { - needNewAlphaTile = true; - } } } else @@ -158,10 +154,10 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst } if (needNewAlphaTile) { - uint _306 = atomic_fetch_add_explicit((device atomic_uint*)&_303.iIndirectDrawParams[4], 1u, memory_order_relaxed); - uint drawBatchAlphaTileIndex = _306; - _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex; - _310.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 1u] = uint(clipAlphaTileIndex); + uint _305 = atomic_fetch_add_explicit((device atomic_uint*)&_302.iIndirectDrawParams[4], 1u, memory_order_relaxed); + uint drawBatchAlphaTileIndex = _305; + _309.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 0u] = drawTileIndex; + _309.iAlphaTiles[(drawBatchAlphaTileIndex * 2u) + 1u] = uint(clipAlphaTileIndex); drawAlphaTileIndex = int(drawBatchAlphaTileIndex) + uFirstAlphaTileIndex; } _175.iDrawTiles[(drawTileIndex * 4u) + 2u] = (uint(drawAlphaTileIndex) & 16777215u) | (uint(drawBackdropDelta) << uint(24)); @@ -170,12 +166,12 @@ kernel void main0(constant int& uColumnCount [[buffer(0)]], constant int& uFirst int tileMapIndex = (tileCoord_1.y * uFramebufferTileSize.x) + tileCoord_1.x; if ((zWrite && (drawTileBackdrop != 0)) && (drawAlphaTileIndex < 0)) { - int _386 = atomic_fetch_max_explicit((device atomic_int*)&_381.iZBuffer[tileMapIndex], int(drawTileIndex), memory_order_relaxed); + int _385 = atomic_fetch_max_explicit((device atomic_int*)&_380.iZBuffer[tileMapIndex], int(drawTileIndex), memory_order_relaxed); } if ((drawTileBackdrop != 0) || (drawAlphaTileIndex >= 0)) { - int _403 = atomic_exchange_explicit((device atomic_int*)&_398.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed); - int nextTileIndex = _403; + int _402 = atomic_exchange_explicit((device atomic_int*)&_397.iFirstTileMap[tileMapIndex], int(drawTileIndex), memory_order_relaxed); + int nextTileIndex = _402; _175.iDrawTiles[(drawTileIndex * 4u) + 0u] = uint(nextTileIndex); } currentBackdrop += drawBackdropDelta; diff --git a/shaders/d3d11/propagate.cs.glsl b/shaders/d3d11/propagate.cs.glsl index 6dd00060..543440f3 100644 --- a/shaders/d3d11/propagate.cs.glsl +++ b/shaders/d3d11/propagate.cs.glsl @@ -148,11 +148,6 @@ void main() { clipTileRect, clipTileCoord); -/* - clipAlphaTileIndex = - int(iClipTiles[clipTileIndex * 4 + - TILE_FIELD_BACKDROP_ALPHA_TILE_ID] << 8) >> 8; - */ int thisClipAlphaTileIndex = int(iClipTiles[clipTileIndex * 4 + TILE_FIELD_BACKDROP_ALPHA_TILE_ID] << 8) >> 8; @@ -166,8 +161,8 @@ void main() { needNewAlphaTile = true; } else { if (drawTileBackdrop != 0) { - // This is a solid draw tile, but there's a clip applied. Replace it with an - // alpha tile pointing directly to the clip mask. + // This is a solid draw tile, but there's a clip applied. Replace it + // with an alpha tile pointing directly to the clip mask. drawAlphaTileIndex = thisClipAlphaTileIndex; clipAlphaTileIndex = -1; needNewAlphaTile = false; @@ -184,8 +179,6 @@ void main() { // This is a blank clip tile. Cull the draw tile entirely. drawTileBackdrop = 0; needNewAlphaTile = false; - } else { - needNewAlphaTile = true; } } } else { diff --git a/svg/src/lib.rs b/svg/src/lib.rs index 1a732a30..79724916 100644 --- a/svg/src/lib.rs +++ b/svg/src/lib.rs @@ -40,7 +40,7 @@ const HAIRLINE_STROKE_WIDTH: f32 = 0.0333; pub struct SVGScene { pub scene: Scene, pub result_flags: BuildResultFlags, - pub clip_paths: HashMap, + pub clip_paths: HashMap, gradients: HashMap, } @@ -110,13 +110,12 @@ impl SVGScene { } if let Some(ref clip_path_name) = group.clip_path { - if let Some(clip_path_id) = self.clip_paths.get(clip_path_name) { - // TODO(pcwalton): Combine multiple clip paths if there's already one. - if state.clip_path.is_some() { - self.result_flags - .insert(BuildResultFlags::UNSUPPORTED_MULTIPLE_CLIP_PATHS); - } - state.clip_path = Some(*clip_path_id); + if let Some(clip_outline) = self.clip_paths.get(clip_path_name) { + let mut clip_path = ClipPath::new((*clip_outline).clone()); + clip_path.set_clip_path(state.clip_path); + clip_path.set_name(format!("ClipPath({})", clip_path_name)); + let clip_path_id = self.scene.push_clip_path(clip_path); + state.clip_path = Some(clip_path_id); } } @@ -187,13 +186,7 @@ impl SVGScene { self.process_node(&kid, &state, &mut clip_outline); } - if let Some(clip_outline) = clip_outline { - // FIXME(pcwalton): Is the winding fill rule correct to use? - let mut clip_path = ClipPath::new(clip_outline); - clip_path.set_name(format!("ClipPath({})", node.id())); - let clip_path_id = self.scene.push_clip_path(clip_path); - self.clip_paths.insert(node.id().to_owned(), clip_path_id); - } + self.clip_paths.insert(node.id().to_owned(), clip_outline.unwrap()); } NodeKind::Defs => { // FIXME(pcwalton): This is wrong.