vello: add AaConfig::SparseMsaa16, backed by the CPU sparse strips The CPU sparse strips (the port of the Skia MSAA sparse strips) stand in for the GPU flatten, path count and path tiling stages. The GPU still runs the draw, clip, binning and tile_alloc stages. New GPU stages: strip_scatter writes the mask tiles and backdrop deltas into the allocated tiles, backdrop_dyn is reused for the prefix sum, coarse_sparse emits the new CMD_MASK, and fine_sparse_msaa16 resolves the 16-bit sample masks. SparseMsaa16 needs the GPU shaders: with use_cpu, rendering returns Error::SparseMsaa16RequiresGpu. with_winit can cycle to the new mode. Tests: end-to-end comparisons against the Area and Msaa16 pipelines on 12 test scenes, and per-pixel comparisons of the Skia rasterizer against the exact reference. Known issue: longpathdash_butt renders one spurious solid tile.
diff --git a/research/examples/with_winit/src/lib.rs b/research/examples/with_winit/src/lib.rs index f795db8..92f0925 100644 --- a/research/examples/with_winit/src/lib.rs +++ b/research/examples/with_winit/src/lib.rs
@@ -110,7 +110,12 @@ #[cfg(not(target_os = "android"))] // TODO: Make this set configurable through the command line // Alternatively, load anti-aliasing shaders on demand/asynchronously -const AA_CONFIGS: [AaConfig; 3] = [AaConfig::Area, AaConfig::Msaa8, AaConfig::Msaa16]; +const AA_CONFIGS: [AaConfig; 4] = [ + AaConfig::Area, + AaConfig::Msaa8, + AaConfig::Msaa16, + AaConfig::SparseMsaa16, +]; #[cfg(target_os = "android")] // Hard code to only one on Android whilst we are working on startup speed @@ -543,6 +548,12 @@ .or(scene_params.base_color) .unwrap_or(palette::css::BLACK); let antialiasing_method = AA_CONFIGS[self.aa_config_ix as usize]; + // The CPU sparse strips can't be combined with `--use-cpu`; use area instead. + #[cfg(not(target_arch = "wasm32"))] + let antialiasing_method = match antialiasing_method { + AaConfig::SparseMsaa16 if self.use_cpu => AaConfig::Area, + method => method, + }; let render_params = vello::RenderParams { base_color, width,
diff --git a/research/examples/with_winit/src/stats.rs b/research/examples/with_winit/src/stats.rs index 4bcf219..ca19e96 100644 --- a/research/examples/with_winit/src/stats.rs +++ b/research/examples/with_winit/src/stats.rs
@@ -64,6 +64,7 @@ AaConfig::Area => "Analytic Area", AaConfig::Msaa16 => "16xMSAA", AaConfig::Msaa8 => "8xMSAA", + AaConfig::SparseMsaa16 => "16xMSAA (CPU sparse strips)", } ), format!("Resolution: {viewport_width}x{viewport_height}"),
diff --git a/research/vello_research/src/lib.rs b/research/vello_research/src/lib.rs index 0b460be..2e1fbca 100644 --- a/research/vello_research/src/lib.rs +++ b/research/vello_research/src/lib.rs
@@ -192,6 +192,14 @@ /// /// Can only be used if [enabled][AaSupport::msaa16] for the `Renderer`. Msaa16, + /// CPU sparse strips with 16-sample (MSAA16) coverage masks. + /// + /// This is a temporary stand-in for GPU sparse strips: paths are rasterized on the CPU + /// into per-pixel sample masks, which the GPU coarse and fine stages then consume. + /// Requires `RendererOptions::use_cpu` to be `false`. + /// + /// Can only be used if [enabled][AaSupport::sparse_msaa16] for the `Renderer`. + SparseMsaa16, } /// Represents the set of anti-aliasing configurations to enable during pipeline creation. @@ -209,6 +217,8 @@ pub msaa8: bool, /// Support [`AaConfig::Msaa16`]. pub msaa16: bool, + /// Support [`AaConfig::SparseMsaa16`]. + pub sparse_msaa16: bool, } impl AaSupport { @@ -220,6 +230,7 @@ area: true, msaa8: true, msaa16: true, + sparse_msaa16: true, } } @@ -231,6 +242,7 @@ area: true, msaa8: false, msaa16: false, + sparse_msaa16: false, } } } @@ -241,12 +253,14 @@ area: false, msaa8: false, msaa16: false, + sparse_msaa16: false, }; for config in iter { match config { AaConfig::Area => result.area = true, AaConfig::Msaa8 => result.msaa8 = true, AaConfig::Msaa16 => result.msaa16 = true, + AaConfig::SparseMsaa16 => result.sparse_msaa16 = true, } } result @@ -295,6 +309,13 @@ #[error("wgpu Error from scope")] WgpuErrorFromScope(#[from] wgpu::Error), + /// [`AaConfig::SparseMsaa16`] was used with a `Renderer` created with + /// `RendererOptions::use_cpu`. The sparse strips already run on the CPU, and their + /// output can only be consumed by the GPU shaders. + #[cfg(feature = "wgpu")] + #[error("`AaConfig::SparseMsaa16` requires `RendererOptions::use_cpu` to be false")] + SparseMsaa16RequiresGpu, + /// Failed to create [`GpuProfiler`]. /// See [`wgpu_profiler::CreationError`] for more information. #[cfg(feature = "wgpu-profiler")] @@ -481,6 +502,7 @@ texture: &TextureView, params: &RenderParams, ) -> Result<()> { + self.check_params(params)?; let (recording, target) = render::render_full( scene, &mut self.resolver, @@ -516,6 +538,16 @@ Ok(()) } + /// Checks that `params` can be rendered by this renderer. + fn check_params(&self, params: &RenderParams) -> Result<()> { + // With `use_cpu`, the sparse strip output would have to flow into the CPU shaders, + // which is not supported. + if params.antialiasing_method == AaConfig::SparseMsaa16 && self.shaders.pathtag_is_cpu { + return Err(Error::SparseMsaa16RequiresGpu); + } + Ok(()) + } + /// Overwrite `image` with `texture`. /// /// Most users should prefer [`register_texture`](Self::register_texture), which @@ -724,6 +756,7 @@ texture: &TextureView, params: &RenderParams, ) -> Result<RenderResult> { + self.check_params(params)?; let mut render = Render::new(); let encoding = scene.encoding(); // TODO: turn this on; the download feature interacts with CPU dispatch.
diff --git a/research/vello_research/src/render.rs b/research/vello_research/src/render.rs index 8c901e3..54387a1 100644 --- a/research/vello_research/src/render.rs +++ b/research/vello_research/src/render.rs
@@ -55,6 +55,24 @@ out_image: ImageProxy, } +/// The resolved scene, and the resources shared by all coarse pipelines. +struct CoarseSetup { + recording: Recording, + /// The packed scene, which is never empty. + packed: Vec<u8>, + cpu_config: vello_encoding::RenderConfig, + gradient_image: ResourceProxy, + image_atlas: ImageProxy, + config_buf: ResourceProxy, +} + +/// Buffers made by `Render::record_draw_to_tile_alloc`, used by the later stages. +struct DrawTileBuffers { + draw_monoid_buf: ResourceProxy, + bin_header_buf: ResourceProxy, + path_buf: ResourceProxy, +} + /// A collection of internal buffers that are used for debug visualization when the /// `debug_layers` feature is enabled. The contents of these buffers remain GPU resident /// and must be freed directly by the caller. @@ -141,6 +159,26 @@ params: &RenderParams, robust: bool, ) -> Recording { + let setup = Self::prepare_coarse(encoding, resolver, persistent_image_atlas, params); + if params.antialiasing_method == AaConfig::SparseMsaa16 { + #[cfg(feature = "wgpu")] + { + return self.render_encoding_coarse_sparse(setup, shaders, params, robust); + } + #[cfg(not(feature = "wgpu"))] + panic!("`AaConfig::SparseMsaa16` requires the `wgpu` feature"); + } + self.render_encoding_coarse_gpu(setup, shaders, params, robust) + } + + /// Resolve the encoding, and prepare the resources shared by all coarse pipelines: + /// the gradient ramps, the image atlas and the render config. + fn prepare_coarse( + encoding: &Encoding, + resolver: &mut Resolver, + persistent_image_atlas: &mut Option<ImageProxy>, + params: &RenderParams, + ) -> CoarseSetup { use vello_encoding::RenderConfig; let mut recording = Recording::default(); let mut packed = vec![]; @@ -216,9 +254,6 @@ params.height ); } - let buffer_sizes = &cpu_config.buffer_sizes; - let wg_counts = &cpu_config.workgroup_counts; - if packed.is_empty() { // HACK: wgpu doesn't allow empty buffers, so we make sure that the scene buffer we upload // can contain at least one array item. @@ -226,10 +261,38 @@ // is zero. packed.resize(size_of::<u32>(), u8::MAX); } - let scene_buf = ResourceProxy::Buffer(recording.upload("vello.scene", packed)); let config_buf = ResourceProxy::Buffer( recording.upload_uniform("vello.config", bytemuck::bytes_of(&cpu_config.gpu)), ); + CoarseSetup { + recording, + packed, + cpu_config, + gradient_image, + image_atlas, + config_buf, + } + } + + /// The GPU coarse pipeline, used for the area and MSAA modes. + fn render_encoding_coarse_gpu( + &mut self, + setup: CoarseSetup, + shaders: &FullShaders, + params: &RenderParams, + robust: bool, + ) -> Recording { + let CoarseSetup { + mut recording, + packed, + cpu_config, + gradient_image, + image_atlas, + config_buf, + } = setup; + let buffer_sizes = &cpu_config.buffer_sizes; + let wg_counts = &cpu_config.workgroup_counts; + let scene_buf = ResourceProxy::Buffer(recording.upload("vello.scene", packed)); let info_bin_data_buf = ResourceProxy::new_buf( buffer_sizes.bin_data.size_in_bytes() as u64, "vello.info_bin_data_buf", @@ -326,115 +389,21 @@ lines_buf, ], ); - let draw_reduced_buf = ResourceProxy::new_buf( - buffer_sizes.draw_reduced.size_in_bytes().into(), - "vello.draw_reduced_buf", + let DrawTileBuffers { + draw_monoid_buf, + bin_header_buf, + path_buf, + } = Self::record_draw_to_tile_alloc( + &mut recording, + shaders, + &cpu_config, + config_buf, + scene_buf, + path_bbox_buf, + info_bin_data_buf, + bump_buf, + tile_buf, ); - recording.dispatch( - shaders.draw_reduce, - wg_counts.draw_reduce, - [config_buf, scene_buf, draw_reduced_buf], - ); - let draw_monoid_buf = ResourceProxy::new_buf( - buffer_sizes.draw_monoids.size_in_bytes().into(), - "vello.draw_monoid_buf", - ); - let clip_inp_buf = ResourceProxy::new_buf( - buffer_sizes.clip_inps.size_in_bytes().into(), - "vello.clip_inp_buf", - ); - recording.dispatch( - shaders.draw_leaf, - wg_counts.draw_leaf, - [ - config_buf, - scene_buf, - draw_reduced_buf, - path_bbox_buf, - draw_monoid_buf, - info_bin_data_buf, - clip_inp_buf, - ], - ); - recording.free_resource(draw_reduced_buf); - let clip_el_buf = ResourceProxy::new_buf( - buffer_sizes.clip_els.size_in_bytes().into(), - "vello.clip_el_buf", - ); - let clip_bic_buf = ResourceProxy::new_buf( - buffer_sizes.clip_bics.size_in_bytes().into(), - "vello.clip_bic_buf", - ); - if wg_counts.clip_reduce.0 > 0 { - recording.dispatch( - shaders.clip_reduce, - wg_counts.clip_reduce, - [clip_inp_buf, path_bbox_buf, clip_bic_buf, clip_el_buf], - ); - } - let clip_bbox_buf = ResourceProxy::new_buf( - buffer_sizes.clip_bboxes.size_in_bytes().into(), - "vello.clip_bbox_buf", - ); - if wg_counts.clip_leaf.0 > 0 { - recording.dispatch( - shaders.clip_leaf, - wg_counts.clip_leaf, - [ - config_buf, - clip_inp_buf, - path_bbox_buf, - clip_bic_buf, - clip_el_buf, - draw_monoid_buf, - clip_bbox_buf, - ], - ); - } - recording.free_resource(clip_inp_buf); - recording.free_resource(clip_bic_buf); - recording.free_resource(clip_el_buf); - let draw_bbox_buf = ResourceProxy::new_buf( - buffer_sizes.draw_bboxes.size_in_bytes().into(), - "vello.draw_bbox_buf", - ); - let bin_header_buf = ResourceProxy::new_buf( - buffer_sizes.bin_headers.size_in_bytes().into(), - "vello.bin_header_buf", - ); - recording.dispatch( - shaders.binning, - wg_counts.binning, - [ - config_buf, - draw_monoid_buf, - path_bbox_buf, - clip_bbox_buf, - draw_bbox_buf, - bump_buf, - info_bin_data_buf, - bin_header_buf, - ], - ); - recording.free_resource(draw_monoid_buf); - recording.free_resource(clip_bbox_buf); - // Note: this only needs to be rounded up because of the workaround to store the tile_offset - // in storage rather than workgroup memory. - let path_buf = - ResourceProxy::new_buf(buffer_sizes.paths.size_in_bytes().into(), "vello.path_buf"); - recording.dispatch( - shaders.tile_alloc, - wg_counts.tile_alloc, - [ - config_buf, - scene_buf, - draw_bbox_buf, - bump_buf, - path_buf, - tile_buf, - ], - ); - recording.free_resource(draw_bbox_buf); recording.free_resource(tagmonoid_buf); let indirect_count_buf = BufferProxy::new( buffer_sizes.indirect_count.size_in_bytes().into(), @@ -556,6 +525,340 @@ recording } + /// Records the draw object, clip, binning and tile allocation stages, which all coarse + /// pipelines share. + fn record_draw_to_tile_alloc( + recording: &mut Recording, + shaders: &FullShaders, + cpu_config: &vello_encoding::RenderConfig, + config_buf: ResourceProxy, + scene_buf: ResourceProxy, + path_bbox_buf: ResourceProxy, + info_bin_data_buf: ResourceProxy, + bump_buf: ResourceProxy, + tile_buf: ResourceProxy, + ) -> DrawTileBuffers { + let buffer_sizes = &cpu_config.buffer_sizes; + let wg_counts = &cpu_config.workgroup_counts; + let draw_reduced_buf = ResourceProxy::new_buf( + buffer_sizes.draw_reduced.size_in_bytes().into(), + "vello.draw_reduced_buf", + ); + recording.dispatch( + shaders.draw_reduce, + wg_counts.draw_reduce, + [config_buf, scene_buf, draw_reduced_buf], + ); + let draw_monoid_buf = ResourceProxy::new_buf( + buffer_sizes.draw_monoids.size_in_bytes().into(), + "vello.draw_monoid_buf", + ); + let clip_inp_buf = ResourceProxy::new_buf( + buffer_sizes.clip_inps.size_in_bytes().into(), + "vello.clip_inp_buf", + ); + recording.dispatch( + shaders.draw_leaf, + wg_counts.draw_leaf, + [ + config_buf, + scene_buf, + draw_reduced_buf, + path_bbox_buf, + draw_monoid_buf, + info_bin_data_buf, + clip_inp_buf, + ], + ); + recording.free_resource(draw_reduced_buf); + let clip_el_buf = ResourceProxy::new_buf( + buffer_sizes.clip_els.size_in_bytes().into(), + "vello.clip_el_buf", + ); + let clip_bic_buf = ResourceProxy::new_buf( + buffer_sizes.clip_bics.size_in_bytes().into(), + "vello.clip_bic_buf", + ); + if wg_counts.clip_reduce.0 > 0 { + recording.dispatch( + shaders.clip_reduce, + wg_counts.clip_reduce, + [clip_inp_buf, path_bbox_buf, clip_bic_buf, clip_el_buf], + ); + } + let clip_bbox_buf = ResourceProxy::new_buf( + buffer_sizes.clip_bboxes.size_in_bytes().into(), + "vello.clip_bbox_buf", + ); + if wg_counts.clip_leaf.0 > 0 { + recording.dispatch( + shaders.clip_leaf, + wg_counts.clip_leaf, + [ + config_buf, + clip_inp_buf, + path_bbox_buf, + clip_bic_buf, + clip_el_buf, + draw_monoid_buf, + clip_bbox_buf, + ], + ); + } + recording.free_resource(clip_inp_buf); + recording.free_resource(clip_bic_buf); + recording.free_resource(clip_el_buf); + let draw_bbox_buf = ResourceProxy::new_buf( + buffer_sizes.draw_bboxes.size_in_bytes().into(), + "vello.draw_bbox_buf", + ); + let bin_header_buf = ResourceProxy::new_buf( + buffer_sizes.bin_headers.size_in_bytes().into(), + "vello.bin_header_buf", + ); + recording.dispatch( + shaders.binning, + wg_counts.binning, + [ + config_buf, + draw_monoid_buf, + path_bbox_buf, + clip_bbox_buf, + draw_bbox_buf, + bump_buf, + info_bin_data_buf, + bin_header_buf, + ], + ); + recording.free_resource(draw_monoid_buf); + recording.free_resource(clip_bbox_buf); + // Note: this only needs to be rounded up because of the workaround to store the tile_offset + // in storage rather than workgroup memory. + let path_buf = + ResourceProxy::new_buf(buffer_sizes.paths.size_in_bytes().into(), "vello.path_buf"); + recording.dispatch( + shaders.tile_alloc, + wg_counts.tile_alloc, + [ + config_buf, + scene_buf, + draw_bbox_buf, + bump_buf, + path_buf, + tile_buf, + ], + ); + recording.free_resource(draw_bbox_buf); + DrawTileBuffers { + draw_monoid_buf, + bin_header_buf, + path_buf, + } + } + + /// The coarse pipeline for [`AaConfig::SparseMsaa16`]. + /// + /// The sparse strips are made on the CPU, which replaces the GPU path tag, bbox and + /// flatten stages, as well as path counting and tiling. The GPU still runs the draw + /// object, clip, binning and tile allocation stages; then `strip_scatter` writes the strip + /// records into the allocated tiles, `backdrop_dyn` prefix-sums their backdrop deltas, + /// and `coarse_sparse` emits `CMD_MASK` and `CMD_SOLID` commands for fine. + #[cfg(feature = "wgpu")] + fn render_encoding_coarse_sparse( + &mut self, + setup: CoarseSetup, + shaders: &FullShaders, + params: &RenderParams, + robust: bool, + ) -> Recording { + use vello_encoding::PathBbox; + use vello_shaders::cpu::sparse_strips::{ + MASK_WORDS_PER_TILE, StripRecord, render_sparse_strips, + }; + + // The workgroup size of `strip_scatter`. + const SCATTER_WG_SIZE: u32 = 256; + // The maximum number of `strip_scatter` workgroups per dispatch row. + const SCATTER_MAX_WGS_X: u32 = 32768; + + assert!( + !shaders.pathtag_is_cpu, + "`AaConfig::SparseMsaa16` is not supported with `RendererOptions::use_cpu`" + ); + let not_configured = "shaders not configured to support AA mode: sparse_msaa16"; + let strip_scatter = shaders.strip_scatter.expect(not_configured); + let coarse_sparse = shaders.coarse_sparse.expect(not_configured); + + let CoarseSetup { + mut recording, + packed, + cpu_config, + gradient_image, + image_atlas, + config_buf, + } = setup; + let buffer_sizes = &cpu_config.buffer_sizes; + let wg_counts = &cpu_config.workgroup_counts; + + // Make the sparse strips on the CPU. (`bytemuck::pod_collect_to_vec` would need + // bytemuck's `extern_crate_alloc` feature.) + let scene: Vec<u32> = packed + .chunks_exact(4) + .map(|w| u32::from_ne_bytes([w[0], w[1], w[2], w[3]])) + .collect(); + let strips = render_sparse_strips(&cpu_config.gpu, &scene); + let n_records = u32::try_from(strips.records.len()).expect("too many sparse strip records"); + // wgpu doesn't allow empty buffers, so pad empty uploads. The padding is never read: + // `n_records` is the true count, and no path or mask block refers to the padding. + let mut path_bboxes = strips.path_bboxes; + if path_bboxes.is_empty() { + path_bboxes.push(PathBbox::default()); + } + let mut records = strips.records; + if records.is_empty() { + records.push(StripRecord::default()); + } + let mut masks = strips.masks; + if masks.is_empty() { + masks.resize(MASK_WORDS_PER_TILE, 0); + } + + let scene_buf = ResourceProxy::Buffer(recording.upload("vello.scene", packed)); + let path_bbox_data: &[u8] = bytemuck::cast_slice(&path_bboxes); + let path_bbox_buf = + ResourceProxy::Buffer(recording.upload("vello.path_bbox_buf", path_bbox_data)); + let record_data: &[u8] = bytemuck::cast_slice(&records); + let records_buf = + ResourceProxy::Buffer(recording.upload("vello.strip_records", record_data)); + let mask_data: &[u8] = bytemuck::cast_slice(&masks); + let masks_buf = ResourceProxy::Buffer(recording.upload("vello.strip_masks", mask_data)); + let info_bin_data_buf = ResourceProxy::new_buf( + buffer_sizes.bin_data.size_in_bytes().into(), + "vello.info_bin_data_buf", + ); + let tile_buf = + ResourceProxy::new_buf(buffer_sizes.tiles.size_in_bytes().into(), "vello.tile_buf"); + let ptcl_buf = + ResourceProxy::new_buf(buffer_sizes.ptcl.size_in_bytes().into(), "vello.ptcl_buf"); + let bump_buf = BufferProxy::new( + buffer_sizes.bump_alloc.size_in_bytes().into(), + "vello.bump_buf", + ); + recording.clear_all(bump_buf); + let bump_buf = ResourceProxy::Buffer(bump_buf); + + let DrawTileBuffers { + draw_monoid_buf, + bin_header_buf, + path_buf, + } = Self::record_draw_to_tile_alloc( + &mut recording, + shaders, + &cpu_config, + config_buf, + scene_buf, + path_bbox_buf, + info_bin_data_buf, + bump_buf, + tile_buf, + ); + recording.free_resource(path_bbox_buf); + + // One invocation per record. The dispatch is split over y when it would need more + // than `SCATTER_MAX_WGS_X` workgroups, and is empty if there are no records. + let n_scatter_wgs = n_records.div_ceil(SCATTER_WG_SIZE); + let scatter_wgs_x = n_scatter_wgs.min(SCATTER_MAX_WGS_X); + let scatter_wgs_y = n_scatter_wgs.div_ceil(SCATTER_MAX_WGS_X); + if n_records != 0 { + // Matches `StripScatterConfig` in `strip_scatter.wgsl`. + let scatter_config: [u32; 4] = [ + n_records, + cpu_config.gpu.layout.n_draw_objects, + scatter_wgs_x * SCATTER_WG_SIZE, + 0, + ]; + let scatter_config_buf = ResourceProxy::Buffer(recording.upload_uniform( + "vello.strip_scatter_config", + bytemuck::bytes_of(&scatter_config), + )); + recording.dispatch( + strip_scatter, + (scatter_wgs_x, scatter_wgs_y, 1), + [ + scatter_config_buf, + records_buf, + path_buf, + bump_buf, + tile_buf, + ], + ); + recording.free_resource(scatter_config_buf); + } + recording.free_resource(records_buf); + recording.dispatch( + shaders.backdrop, + wg_counts.backdrop, + [config_buf, bump_buf, path_buf, tile_buf], + ); + recording.dispatch( + coarse_sparse, + wg_counts.coarse, + [ + config_buf, + scene_buf, + draw_monoid_buf, + bin_header_buf, + info_bin_data_buf, + path_buf, + tile_buf, + bump_buf, + ptcl_buf, + ], + ); + // Still needed, as it writes the `ptcl[0] = ~0` failure marker which fine checks. The + // indirect dispatch size it computes (for path tiling) is unused. + let indirect_count_buf = BufferProxy::new( + buffer_sizes.indirect_count.size_in_bytes().into(), + "vello.indirect_count", + ); + recording.dispatch( + shaders.path_tiling_setup, + wg_counts.path_tiling_setup, + [bump_buf, indirect_count_buf.into(), ptcl_buf], + ); + recording.free_buffer(indirect_count_buf); + recording.free_resource(scene_buf); + recording.free_resource(draw_monoid_buf); + recording.free_resource(bin_header_buf); + recording.free_resource(path_buf); + let out_image = ImageProxy::new(params.width, params.height, ImageFormat::Rgba8); + let blend_spill_buf = BufferProxy::new( + buffer_sizes.blend_spill.size_in_bytes().into(), + "vello.blend_spill", + ); + self.fine_wg_count = Some(wg_counts.fine); + self.fine_resources = Some(FineResources { + aa_config: params.antialiasing_method, + config_buf, + bump_buf, + tile_buf, + // Fine reads the masks through the segments binding. + segments_buf: masks_buf, + ptcl_buf, + gradient_image, + info_bin_data_buf, + blend_spill_buf: ResourceProxy::Buffer(blend_spill_buf), + image_atlas: ResourceProxy::Image(image_atlas), + out_image, + }); + if robust { + recording.download(*bump_buf.as_buf().unwrap()); + } + recording.free_resource(bump_buf); + // There are no flattened lines to capture for the debug layers. + recording + } + /// Run fine rasterization assuming the coarse phase succeeded. pub fn record_fine(&mut self, shaders: &FullShaders, recording: &mut Recording) { let fine_wg_count = self.fine_wg_count.take().unwrap(); @@ -579,6 +882,25 @@ ], ); } + AaConfig::SparseMsaa16 => { + // The CPU masks are bound in place of the segments. + recording.dispatch( + shaders + .fine_sparse_msaa16 + .expect("shaders not configured to support AA mode: sparse_msaa16"), + fine_wg_count, + [ + fine.config_buf, + fine.segments_buf, + fine.ptcl_buf, + fine.info_bin_data_buf, + fine.blend_spill_buf, + ResourceProxy::Image(fine.out_image), + fine.gradient_image, + fine.image_atlas, + ], + ); + } _ => { if self.mask_buf.is_none() { let mask_lut = match fine.aa_config {
diff --git a/research/vello_research/src/shaders.rs b/research/vello_research/src/shaders.rs index 6311637..df81aa6 100644 --- a/research/vello_research/src/shaders.rs +++ b/research/vello_research/src/shaders.rs
@@ -39,6 +39,10 @@ pub fine_area: Option<ShaderId>, pub fine_msaa8: Option<ShaderId>, pub fine_msaa16: Option<ShaderId>, + // Sparse strips (`AaConfig::SparseMsaa16`). Only built if supported and not `use_cpu`. + pub strip_scatter: Option<ShaderId>, + pub coarse_sparse: Option<ShaderId>, + pub fine_sparse_msaa16: Option<ShaderId>, // 2-level dispatch works for CPU pathtag scan even for large // inputs, 3-level is not yet implemented. pub pathtag_is_cpu: bool, @@ -245,6 +249,46 @@ } else { None }; + // The sparse strips are produced on the CPU and consumed by these GPU-only shaders, + // so they can't be combined with `use_cpu`. + let sparse = aa_support.sparse_msaa16 && !options.use_cpu; + let strip_scatter = if sparse { + Some(add_shader!( + strip_scatter, + [Uniform, BufReadOnly, BufReadOnly, Buffer, Buffer], + CpuShaderType::Missing + )) + } else { + None + }; + let coarse_sparse = if sparse { + Some(add_shader!( + coarse_sparse, + [ + Uniform, + BufReadOnly, + BufReadOnly, + BufReadOnly, + BufReadOnly, + BufReadOnly, + Buffer, + Buffer, + Buffer, + ], + CpuShaderType::Missing + )) + } else { + None + }; + let fine_sparse_msaa16 = if sparse { + Some(add_shader!( + fine_sparse_msaa16, + fine_resources[..fine_resources.len() - 1], + CpuShaderType::Missing + )) + } else { + None + }; Ok(FullShaders { pathtag_reduce, @@ -269,6 +313,9 @@ fine_area, fine_msaa8, fine_msaa16, + strip_scatter, + coarse_sparse, + fine_sparse_msaa16, pathtag_is_cpu: options.use_cpu, }) }
diff --git a/research/vello_research_tests/tests/sparse_strips.rs b/research/vello_research_tests/tests/sparse_strips.rs new file mode 100644 index 0000000..524c821 --- /dev/null +++ b/research/vello_research_tests/tests/sparse_strips.rs
@@ -0,0 +1,195 @@ +// Copyright 2026 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +//! End-to-end tests for [`AaConfig::SparseMsaa16`]: the CPU sparse strips (a port of Skia's +//! MSAA sparse strips) feeding the GPU draw, clip, coarse and fine stages. +//! +//! Each scene is rendered with the sparse strips and with the GPU `Area` and `Msaa16` +//! pipelines. The pipelines only differ in how they anti-alias edges, so the sparse strips +//! must be about as close to `Area` as the GPU's own 16x MSAA is. Missing tiles, wrong +//! winding or leaking clips show up as much larger differences. + +use std::path::Path; + +use nv_flip::{FlipImageRgb8, FlipPool}; +use scenes::{ExampleScene, test_scenes}; +use vello::peniko::ImageData; +use vello::{AaConfig, Scene}; +use vello_research_tests::{ + TestParams, encode_test_scene, render_then_debug_sync, write_png_to_file, +}; + +fn flip_image(image: &ImageData) -> FlipImageRgb8 { + let rgb: Vec<u8> = image + .data + .data() + .chunks_exact(4) + .flat_map(|p| [p[0], p[1], p[2]]) + .collect(); + FlipImageRgb8::with_data(image.width, image.height, &rgb) +} + +/// The nv-flip mean error between two renders, and their largest channel difference. +fn diff(reference: &ImageData, test: &ImageData) -> (f32, u8) { + let error_map = nv_flip::flip( + flip_image(reference), + flip_image(test), + nv_flip::DEFAULT_PIXELS_PER_DEGREE, + ); + let mean = FlipPool::from_image(&error_map).mean(); + let max = reference + .data + .data() + .iter() + .zip(test.data.data()) + .map(|(a, b)| a.abs_diff(*b)) + .max() + .unwrap_or(0); + (mean, max) +} + +fn render(scene: &Scene, params: &mut TestParams, aa: AaConfig) -> ImageData { + params.anti_aliasing = aa; + render_then_debug_sync(scene, params).unwrap() +} + +fn check_scene(test_scene: ExampleScene, name: &str, width: u32, height: u32) { + let mut params = TestParams::new(format!("sparse_strips_{name}"), width, height); + let scene = encode_test_scene(test_scene, &mut params); + let sparse = render(&scene, &mut params, AaConfig::SparseMsaa16); + let area = render(&scene, &mut params, AaConfig::Area); + let msaa16 = render(&scene, &mut params, AaConfig::Msaa16); + + let (sparse_vs_area, sparse_max) = diff(&area, &sparse); + let (msaa16_vs_area, msaa16_max) = diff(&area, &msaa16); + let (sparse_vs_msaa16, _) = diff(&msaa16, &sparse); + println!( + "sparse_strips {name}: flip vs Area: SparseMsaa16 {sparse_vs_area:.5} (max channel \ + diff {sparse_max}), Msaa16 {msaa16_vs_area:.5} (max {msaa16_max}); \ + SparseMsaa16 vs Msaa16 {sparse_vs_msaa16:.5}" + ); + + let limit = 1.5 * msaa16_vs_area + 0.003; + let failed = sparse_vs_area > limit; + // Set `VELLO_SPARSE_STRIPS_DUMP` to write the renders for visual inspection. + if failed || std::env::var_os("VELLO_SPARSE_STRIPS_DUMP").is_some() { + let dir = Path::new(env!("CARGO_MANIFEST_DIR")).join("debug_outputs"); + std::fs::create_dir_all(&dir).unwrap(); + for (image, suffix) in [(&sparse, "sparse"), (&area, "area"), (&msaa16, "msaa16")] { + let path = dir.join(format!("sparse_strips_{name}_{suffix}.png")); + write_png_to_file(¶ms, &path, image, None, false).unwrap(); + } + } + assert!( + !failed, + "{name}: SparseMsaa16 is too far from Area: flip mean {sparse_vs_area} > {limit}. \ + Wrote the renders to debug_outputs/" + ); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_splash() { + check_scene(test_scenes::splash_with_tiger(), "splash", 300, 300); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_funky_paths() { + check_scene(test_scenes::funky_paths(), "funky_paths", 600, 600); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_stroke_styles() { + check_scene(test_scenes::stroke_styles(), "stroke_styles", 600, 425); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_tricky_strokes() { + check_scene(test_scenes::tricky_strokes(), "tricky_strokes", 600, 425); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_fill_types() { + check_scene(test_scenes::fill_types(), "fill_types", 700, 350); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_deep_blend() { + check_scene(test_scenes::deep_blend(), "deep_blend", 200, 200); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_many_clips() { + check_scene(test_scenes::many_clips(), "many_clips", 200, 200); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_clip_test() { + check_scene(test_scenes::clip_test(), "clip_test", 512, 768); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_longpathdash_butt() { + // Thousands of tiny dashes. The sparse strips use Vello's CPU flattener, so also measure how + // far Vello's CPU shaders (`use_cpu`, flatten included) are from the GPU ones, both rendered + // with `Area` by the GPU fine stage. + // + // TODO: The sparse render has one spurious solid tile near the bottom right (pixels about + // 416..432 x 64..80) that the GPU pipelines don't. The flip mean is too coarse to catch it. + let mut params = TestParams::new("sparse_strips_longpathdash_butt_flatten", 440, 80); + let scene = encode_test_scene(test_scenes::longpathdash_butt(), &mut params); + let gpu_flatten = render(&scene, &mut params, AaConfig::Area); + params.use_cpu = true; + let cpu_flatten = render(&scene, &mut params, AaConfig::Area); + let (mean, max) = diff(&gpu_flatten, &cpu_flatten); + println!("sparse_strips longpathdash_butt: Area, CPU vs GPU flatten: {mean:.5} (max {max})"); + check_scene( + test_scenes::longpathdash_butt(), + "longpathdash_butt", + 440, + 80, + ); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_blurred_rounded_rect() { + check_scene( + test_scenes::blurred_rounded_rect(), + "blurred_rounded_rect", + 400, + 400, + ); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_gradient_extend() { + check_scene(test_scenes::gradient_extend(), "gradient_extend", 200, 200); +} + +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_luminance_mask() { + check_scene(test_scenes::luminance_mask(), "luminance_mask", 55, 55); +} + +/// The sparse strips run on the CPU and feed GPU-only shaders, so they can't be combined with +/// `use_cpu`. +#[test] +#[cfg_attr(skip_gpu_tests, ignore)] +fn sparse_strips_reject_use_cpu() { + let mut params = TestParams::new("sparse_strips_reject_use_cpu", 64, 64); + let scene = encode_test_scene(test_scenes::fill_types(), &mut params); + params.anti_aliasing = AaConfig::SparseMsaa16; + params.use_cpu = true; + assert!(render_then_debug_sync(&scene, ¶ms).is_err()); +}
diff --git a/research/vello_shaders/shader/coarse.wgsl b/research/vello_shaders/shader/coarse.wgsl index 4e045c1..308491f 100644 --- a/research/vello_shaders/shader/coarse.wgsl +++ b/research/vello_shaders/shader/coarse.wgsl
@@ -86,6 +86,22 @@ } fn write_path(tile: Tile, tile_ix: u32, draw_flags: u32) { +#ifdef sparse + // Sparse strips (`AaConfig::SparseMsaa16`): `strip_scatter` stores the CPU mask + // block index plus one, and leaves zero in tiles fully covered by the path. + // The fill rule has already been applied by the CPU. + let mask_ix = tile.segment_count_or_ix; + if mask_ix != 0u { + alloc_cmd(2u); + ptcl[cmd_offset] = CMD_MASK; + ptcl[cmd_offset + 1u] = mask_ix - 1u; + cmd_offset += 2u; + } else { + alloc_cmd(1u); + ptcl[cmd_offset] = CMD_SOLID; + cmd_offset += 1u; + } +#else // We overload the "segments" field to store both count (written by // path_count stage) and segment allocation (used by path_tiling and // fine). @@ -107,6 +123,7 @@ ptcl[cmd_offset] = CMD_SOLID; cmd_offset += 1u; } +#endif } fn write_color(color: CmdColor) {
diff --git a/research/vello_shaders/shader/fine.wgsl b/research/vello_shaders/shader/fine.wgsl index 55f615d..bf56e8f 100644 --- a/research/vello_shaders/shader/fine.wgsl +++ b/research/vello_shaders/shader/fine.wgsl
@@ -17,8 +17,15 @@ @group(0) @binding(0) var<uniform> config: Config; +#ifdef sparse +// Sparse strips (`AaConfig::SparseMsaa16`): 16-sample coverage masks written by the +// CPU, 128 words per 16x16 tile. See `vello_shaders::cpu::sparse_strips`. +@group(0) @binding(1) +var<storage> masks: array<u32>; +#else @group(0) @binding(1) var<storage> segments: array<Segment>; +#endif #import blend #import ptcl @@ -995,6 +1002,7 @@ const PIXELS_PER_THREAD = 4u; #ifndef msaa +#ifndef sparse // Analytic area anti-aliasing. // @@ -1058,6 +1066,7 @@ *result = area; } +#endif // sparse #endif // The X size should be 16 / PIXELS_PER_THREAD @@ -1093,6 +1102,22 @@ break; } switch tag { +#ifdef sparse + case CMD_MASK: { + // Pixel (x, y) of the tile is the u16 at index y * 16 + x of the mask + // block, and bit k is sample k. This thread owns x = 4 * local_id.x + i. + let block = ptcl[cmd_ix + 1u]; + let base = block * 128u + local_id.y * 8u + local_id.x * 2u; + let w0 = masks[base]; + let w1 = masks[base + 1u]; + // assumes PIXELS_PER_THREAD == 4 + area[0] = f32(countOneBits(w0 & 0xffffu)) * 0.0625; + area[1] = f32(countOneBits(w0 >> 16u)) * 0.0625; + area[2] = f32(countOneBits(w1 & 0xffffu)) * 0.0625; + area[3] = f32(countOneBits(w1 >> 16u)) * 0.0625; + cmd_ix += 2u; + } +#else case CMD_FILL: { let fill = read_fill(cmd_ix); #ifdef msaa @@ -1102,6 +1127,7 @@ #endif cmd_ix += 4u; } +#endif case CMD_SOLID: { for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { area[i] = 1.0;
diff --git a/research/vello_shaders/shader/permutations b/research/vello_shaders/shader/permutations index b235168..313e88e 100644 --- a/research/vello_shaders/shader/permutations +++ b/research/vello_shaders/shader/permutations
@@ -5,3 +5,7 @@ + fine_area + fine_msaa8: msaa msaa8 + fine_msaa16: msaa msaa16 ++ fine_sparse_msaa16: sparse +coarse ++ coarse ++ coarse_sparse: sparse
diff --git a/research/vello_shaders/shader/shared/ptcl.wgsl b/research/vello_shaders/shader/shared/ptcl.wgsl index 487c349..6342c21 100644 --- a/research/vello_shaders/shader/shared/ptcl.wgsl +++ b/research/vello_shaders/shader/shared/ptcl.wgsl
@@ -23,6 +23,8 @@ const CMD_END_CLIP = 11u; const CMD_JUMP = 12u; const CMD_BLUR_RECT = 13u; +// Sparse strips (`AaConfig::SparseMsaa16`): CPU-provided 16-sample coverage mask. +const CMD_MASK = 14u; // The individual PTCL structs are written here, but read/write is by // hand in the relevant shaders
diff --git a/research/vello_shaders/shader/strip_scatter.wgsl b/research/vello_shaders/shader/strip_scatter.wgsl new file mode 100644 index 0000000..8117fb9 --- /dev/null +++ b/research/vello_shaders/shader/strip_scatter.wgsl
@@ -0,0 +1,95 @@ +// Copyright 2026 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense + +// Scatter the CPU sparse strip records (`AaConfig::SparseMsaa16`) into the tile +// array allocated by `tile_alloc`. +// +// See `vello_shaders::cpu::sparse_strips` for the CPU -> GPU contract. Each record +// is either a mask tile, which sets `segment_count_or_ix = block + 1` (coarse then +// emits `CMD_MASK`), or a backdrop delta, which `backdrop_dyn` (run afterwards) +// prefix-sums into backdrops, so tiles inside solid spans get a backdrop of 1. +// +// The tile rectangle of a path comes from its clip-intersected bbox, so records +// outside it are dropped. A delta left of the rectangle is moved to its left edge, +// which keeps spans that start left of the rectangle and end inside it correct. + +#import bump +#import tile + +struct StripScatterConfig { + n_records: u32, + n_drawobj: u32, + // Invocations per row of the dispatch. The dispatch is split over y when it + // needs more than 65535 workgroups. + row_stride: u32, + _padding: u32, +} + +// Must match `StripRecord` in `vello_shaders::cpu::sparse_strips::sink`. +struct StripRecord { + path_ix: u32, + // tile_x | (tile_y << 16) + xy: u32, + // Mask tile: the mask block index. Backdrop delta: KIND_DELTA | (delta as u16). + payload: u32, +} + +// Same layout as `Tile`, but with atomic fields. +struct AtomicTile { + backdrop: atomic<i32>, + segment_count_or_ix: atomic<u32>, +} + +@group(0) @binding(0) +var<uniform> scatter_config: StripScatterConfig; + +@group(0) @binding(1) +var<storage> records: array<StripRecord>; + +@group(0) @binding(2) +var<storage> paths: array<Path>; + +@group(0) @binding(3) +var<storage, read_write> bump: BumpAllocators; + +@group(0) @binding(4) +var<storage, read_write> tiles: array<AtomicTile>; + +const KIND_DELTA = 0x80000000u; + +@compute @workgroup_size(256) +fn main( + @builtin(global_invocation_id) global_id: vec3<u32>, +) { + // Abort if any of the prior stages failed: the tile allocation can't be trusted. + if atomicLoad(&bump.failed) != 0u { + return; + } + let ix = global_id.y * scatter_config.row_stride + global_id.x; + if ix >= scatter_config.n_records { + return; + } + let record = records[ix]; + if record.path_ix >= scatter_config.n_drawobj { + return; + } + let path = paths[record.path_ix]; + let x = record.xy & 0xffffu; + let y = record.xy >> 16u; + if y < path.bbox.y || y >= path.bbox.w { + return; + } + let row = path.tiles + (y - path.bbox.y) * (path.bbox.z - path.bbox.x); + if (record.payload & KIND_DELTA) == 0u { + if x >= path.bbox.x && x < path.bbox.z { + atomicStore(&tiles[row + x - path.bbox.x].segment_count_or_ix, record.payload + 1u); + } + } else { + let cx = max(x, path.bbox.x); + if cx < path.bbox.z { + // Sign extend the low 16 bits. + let delta = bitcast<i32>(record.payload << 16u) >> 16u; + atomicAdd(&tiles[row + cx - path.bbox.x].backdrop, delta); + } + } +}
diff --git a/research/vello_shaders/src/cpu/sparse_strips/tests.rs b/research/vello_shaders/src/cpu/sparse_strips/tests.rs index 017598e..a994045 100644 --- a/research/vello_shaders/src/cpu/sparse_strips/tests.rs +++ b/research/vello_shaders/src/cpu/sparse_strips/tests.rs
@@ -305,3 +305,139 @@ assert_eq!(sink.records.len(), 3); assert_eq!(sink.records[2], StripRecord::delta(3, 6, 0, 1)); } + +/// Statistics of the per-pixel difference between two decoded outputs over the viewport. +#[derive(Clone, Copy, Debug, Default)] +struct DiffStats { + /// Sum over pixels of |popcount difference|. + popcount: u64, + /// Sum over pixels of the number of differing samples. + samples: u64, + /// Number of differing pixels. + pixels: u64, + /// Largest |popcount difference| of a pixel. + max: u32, +} + +impl DiffStats { + fn of(a: &[u16], b: &[u16], width: u32, height: u32) -> Self { + let stride = (width.div_ceil(TILE_SIZE as u32) * TILE_SIZE as u32) as usize; + let mut stats = Self::default(); + for y in 0..height as usize { + for x in 0..width as usize { + let (a, b) = (a[y * stride + x], b[y * stride + x]); + let d = a.count_ones().abs_diff(b.count_ones()); + stats.popcount += u64::from(d); + stats.samples += u64::from((a ^ b).count_ones()); + stats.pixels += u64::from(a != b); + stats.max = stats.max.max(d); + } + } + stats + } + + fn add(&mut self, other: Self) { + self.popcount += other.popcount; + self.samples += other.samples; + self.pixels += other.pixels; + self.max = self.max.max(other.max); + } +} + +fn perimeter(lines: &[Line]) -> f64 { + lines + .iter() + .map(|l| f64::from(l.p1.x - l.p0.x).hypot(f64::from(l.p1.y - l.p0.y))) + .sum() +} + +/// A random test path: a star polygon for even `i`, else a random (generally self-intersecting) +/// polygon. Both may extend past any edge of a `width` x `height` viewport. +fn random_path(rng: &mut XorShift, i: usize, width: u32, height: u32) -> Vec<Line> { + let (w, h) = (width as f32, height as f32); + let n = 3 + i % 14; + if i.is_multiple_of(2) { + let cx = -0.1 * w + 1.2 * w * rng.next_f32(); + let cy = -0.1 * h + 1.2 * h * rng.next_f32(); + let r = 4.0 + 0.5 * w.min(h) * rng.next_f32(); + star_polygon(rng, cx, cy, r, n).0 + } else { + let pts: Vec<(f32, f32)> = (0..n) + .map(|_| { + ( + -0.15 * w + 1.3 * w * rng.next_f32(), + -0.15 * h + 1.3 * h * rng.next_f32(), + ) + }) + .collect(); + polygon(&pts) + } +} + +/// The LUT quantizes the line offset and slope, so samples very close to an edge may differ +/// from the exact reference. This bounds the summed |popcount difference| of a path relative to +/// its perimeter (which overestimates the edge length inside the viewport). +fn max_popcount_diff(perimeter: f64) -> f64 { + 0.25 * perimeter + 8.0 +} + +#[test] +fn skia_matches_reference() { + let mut rng = XorShift(0x2545_f491_4f6c_dd1d); + // A ragged viewport. + let (w, h) = (200, 150); + let mut total = DiffStats::default(); + let mut total_perimeter = 0.0; + for i in 0..80 { + let lines = random_path(&mut rng, i, w, h); + let perimeter = perimeter(&lines); + for even_odd in [false, true] { + let skia = render(&lines, even_odd, w, h, Backend::Skia); + let reference = render(&lines, even_odd, w, h, Backend::Reference); + let stats = DiffStats::of(&skia, &reference, w, h); + assert!( + stats.popcount as f64 <= max_popcount_diff(perimeter), + "path {i} even_odd {even_odd}: {stats:?}, perimeter {perimeter}: {lines:?}" + ); + total.add(stats); + total_perimeter += perimeter; + } + } + println!( + "skia vs reference, per px of perimeter: |popcount diff| {:.4}, differing samples {:.4}, \ + differing pixels {:.4}; max |popcount diff| of a pixel {}", + total.popcount as f64 / total_perimeter, + total.samples as f64 / total_perimeter, + total.pixels as f64 / total_perimeter, + total.max, + ); +} + +#[test] +fn skia_matches_reference_across_paths() { + // One rasterizer for a whole scene, as in `render_sparse_strips_with`: the winding histogram + // and tile buffers are reused across paths, and must be reset correctly in between. + let mut rng = XorShift(0x0123_4567_89ab_cdef); + let (w, h) = (130, 100); + let paths: Vec<Vec<Line>> = (0..40).map(|i| random_path(&mut rng, i, w, h)).collect(); + let mut skia = StripSink::new(w.div_ceil(TILE_SIZE as u32)); + let mut reference = StripSink::new(w.div_ceil(TILE_SIZE as u32)); + let mut rasterizer = Rasterizer::new(w, h); + for (path_ix, lines) in paths.iter().enumerate() { + let even_odd = path_ix % 3 == 0; + skia.begin_path(path_ix as u32); + rasterizer.rasterize_path(lines, even_odd, &mut skia); + reference.begin_path(path_ix as u32); + rasterize_path_reference(lines, even_odd, w, h, &mut reference); + } + for (path_ix, lines) in paths.iter().enumerate() { + let a = decode(&skia.records, &skia.masks, path_ix as u32, w, h); + let b = decode(&reference.records, &reference.masks, path_ix as u32, w, h); + let stats = DiffStats::of(&a, &b, w, h); + let perimeter = perimeter(lines); + assert!( + stats.popcount as f64 <= max_popcount_diff(perimeter), + "path {path_ix}: {stats:?}, perimeter {perimeter}: {lines:?}" + ); + } +}