prototype gpu merge
diff --git a/Cargo.lock b/Cargo.lock index 82f33ab..079fa00 100644 --- a/Cargo.lock +++ b/Cargo.lock
@@ -3922,6 +3922,8 @@ name = "vello_hybrid_winit" version = "0.0.0" dependencies = [ + "bytemuck", + "futures-intrusive", "pollster", "vello_common", "vello_example_scenes",
diff --git a/sparse_strips/vello_bench/src/data.rs b/sparse_strips/vello_bench/src/data.rs index 989eeec..b860060 100644 --- a/sparse_strips/vello_bench/src/data.rs +++ b/sparse_strips/vello_bench/src/data.rs
@@ -147,6 +147,7 @@ /// Get the alpha buffer and rendered strips. pub fn strips(&self) -> (Vec<u8>, Vec<Strip>) { let mut strip_buf = vec![]; + let mut pmt_buf = vec![]; let mut alpha_buf = vec![]; let lines = self.lines(); let tiles = self.sorted_tiles(); @@ -155,6 +156,7 @@ Level::fallback(), &tiles, &mut strip_buf, + &mut pmt_buf, &mut alpha_buf, Fill::NonZero, None,
diff --git a/sparse_strips/vello_bench/src/strip.rs b/sparse_strips/vello_bench/src/strip.rs index 8690941..f395dec 100644 --- a/sparse_strips/vello_bench/src/strip.rs +++ b/sparse_strips/vello_bench/src/strip.rs
@@ -17,6 +17,7 @@ g.bench_function(format!("{}_{}", $item.name.clone(), $suffix), |b| { let mut strip_buf = vec![]; + let mut pmt_buf = vec![]; let mut alpha_buf = vec![]; b.iter(|| { @@ -27,6 +28,7 @@ $level, &tiles, &mut strip_buf, + &mut pmt_buf, &mut alpha_buf, Fill::NonZero, None,
diff --git a/sparse_strips/vello_common/src/recording.rs b/sparse_strips/vello_common/src/recording.rs index e840e84..481364f 100644 --- a/sparse_strips/vello_common/src/recording.rs +++ b/sparse_strips/vello_common/src/recording.rs
@@ -11,7 +11,7 @@ #[cfg(feature = "text")] use crate::peniko::FontData; use crate::peniko::{BlendMode, Fill}; -use crate::strip::Strip; +use crate::strip::{PreMergeTile, Strip}; use crate::strip_generator::StripStorage; use alloc::vec::Vec; @@ -59,6 +59,10 @@ &self.strip_storage.strips } + pub fn pre_merge_tiles(&self) -> &[PreMergeTile] { + &self.strip_storage.pre_merge_tiles + } + /// Get alphas as slice pub fn alphas(&self) -> &[u8] { &self.strip_storage.alphas @@ -177,8 +181,12 @@ } /// Get cached strips. - pub fn get_cached_strips(&self) -> (&[Strip], &[u8]) { - (self.cached_strips.strips(), self.cached_strips.alphas()) + pub fn get_cached_strips(&self) -> (&[Strip], &[PreMergeTile], &[u8]) { + ( + self.cached_strips.strips(), + self.cached_strips.pre_merge_tiles(), + self.cached_strips.alphas(), + ) } /// Takes cached strip buffers.
diff --git a/sparse_strips/vello_common/src/strip.rs b/sparse_strips/vello_common/src/strip.rs index f2e5b7e..c8f99be 100644 --- a/sparse_strips/vello_common/src/strip.rs +++ b/sparse_strips/vello_common/src/strip.rs
@@ -2,13 +2,17 @@ // SPDX-License-Identifier: Apache-2.0 OR MIT //! Rendering strips. - use crate::flatten::Line; use crate::peniko::Fill; use crate::tile::{Tile, Tiles}; use crate::util::f32_to_u8; use alloc::vec::Vec; +use bytemuck::{Pod, Zeroable}; use fearless_simd::*; +use std::format; +use std::println; +use std::string::String; +use std::vec; /// A strip. #[derive(Debug, Clone, Copy)] @@ -85,17 +89,115 @@ } } +/// Pre-Merge-Tile: +#[repr(C)] +#[derive(Debug, Clone, Copy, Zeroable, Pod)] +pub struct PreMergeTile { + /// The index into the alpha buffer that an "end tile" should write to. + pub alpha_index: u32, + /// Contains tile location information + pub packed_info: u32, + /// The exclusive prefix sum of the signed winding number, uploading it skips scan in shader. + pub scanned_winding: i32, + /// Necessary padding for alignment + pub padding: u32, + /// Line points, adjusted by tile position + pub p0: [f32; 2], + /// Line points, adjusted by tile position + pub p1: [f32; 2], +} + +// MSB LSB +// 31------------------21|20------------------12|11-------------------3|2|1|0| +// | Unused (11) | Seg Start ID (9) | Tile Start ID (9) |F|T|E| +// +// F = FILL_RULE_MASK +// E = IS_END_TILE_MASK +// T = IS_TILE_FIRST_COL_MASK +const IS_END_TILE_MASK: u32 = 1 << 0; +const IS_TILE_FIRST_COL_MASK: u32 = 1 << 1; +const FILL_RULE_MASK: u32 = 1 << 2; +const TILE_START_ID_SHIFT: u32 = 3; +const SEG_START_ID_SHIFT: u32 = 12; +const ID_MASK: u32 = 0x1ff; +const INVALID_ID: u32 = 256; +const BLOCK_DIM_LG: u32 = 8; + +impl PreMergeTile { + pub fn set_is_end_tile(&mut self, is_end: bool) { + if is_end { + self.packed_info |= IS_END_TILE_MASK; + } else { + self.packed_info &= !IS_END_TILE_MASK; + } + } + + pub fn is_end_tile(&self) -> bool { + (self.packed_info & IS_END_TILE_MASK) != 0 + } + + pub fn set_is_tile_first_col(&mut self, is_first: bool) { + if is_first { + self.packed_info |= IS_TILE_FIRST_COL_MASK; + } + } + + pub fn is_tile_first_col(&self) -> bool { + (self.packed_info & IS_TILE_FIRST_COL_MASK) != 0 + } + + pub fn set_fill_rule(& mut self, is_fill_rule_non_zero: bool) { + if is_fill_rule_non_zero { + self.packed_info |= FILL_RULE_MASK; + } + } + + pub fn set_tile_start_id(&mut self, id: u32) { + self.packed_info |= (id & ID_MASK) << TILE_START_ID_SHIFT; + } + + pub fn tile_start_id(&self) -> u32 { + (self.packed_info >> TILE_START_ID_SHIFT) & ID_MASK + } + + pub fn set_seg_start_id(&mut self, id: u32) { + self.packed_info |= (id & ID_MASK) << SEG_START_ID_SHIFT; + } + + pub fn seg_start_id(&self) -> u32 { + (self.packed_info >> SEG_START_ID_SHIFT) & ID_MASK + } +} + /// Render the tiles stored in `tiles` into the strip and alpha buffer. +/// Aliasing threshold is unused pub fn render( - level: Level, + _level: Level, tiles: &Tiles, - strip_buf: &mut Vec<Strip>, - alpha_buf: &mut Vec<u8>, + strip_buf: &mut Vec<Strip>, // Output parameter for strips + pmt_buf: &mut Vec<PreMergeTile>, // Output paramter for pre merge tiles + alpha_buf: &mut Vec<u8>, // Output parameter for alpha bytes fill_rule: Fill, aliasing_threshold: Option<u8>, lines: &[Line], ) { - dispatch!(level, simd => render_impl(simd, tiles, strip_buf, alpha_buf, fill_rule, aliasing_threshold, lines)); + // let mut winding_fine_ref: Vec<[[f32; 4]; 4]> = Vec::new(); + // let mut winding_acc_ref: Vec<[f32; 4]> = Vec::new(); + // let mut winding_coarse_ref: Vec<i32> = Vec::new(); + // let mut winding_fine_comp: Vec<[[f32; 4]; 4]> = Vec::new(); + // let mut winding_acc_comp: Vec<[f32; 4]> = Vec::new(); + // let mut winding_coarse_comp: Vec<i32> = Vec::new(); + + prepare_gpu_inputs(tiles, strip_buf, pmt_buf, alpha_buf, fill_rule, lines); + + // compare_windings( + // &winding_fine_ref, + // &winding_coarse_ref, + // &winding_acc_ref, + // &winding_fine_comp, + // &winding_coarse_comp, + // &winding_acc_comp, + // ); } fn render_impl<S: Simd>( @@ -106,6 +208,9 @@ fill_rule: Fill, aliasing_threshold: Option<u8>, lines: &[Line], + winding_fine_ref: &mut Vec<[[f32; 4]; 4]>, + winding_acc_ref: &mut Vec<[f32; 4]>, + winding_coarse_ref: &mut Vec<i32>, ) { if tiles.is_empty() { return; @@ -155,6 +260,10 @@ // Push out the winding as an alpha mask when we move to the next location (i.e., a tile // without the same location). if !prev_tile.same_loc(&tile) { + //winding_fine_ref.push(location_winding.map(|v| v.val)); + // winding_acc_ref.push(accumulated_winding.val); + // winding_coarse_ref.push(winding_delta); + match fill_rule { Fill::NonZero => { let p1 = f32x4::splat(s, 0.5); @@ -203,7 +312,7 @@ ); } - alpha_buf.extend_from_slice(&u8_vals.val); + //alpha_buf.extend_from_slice(&u8_vals.val); #[expect(clippy::needless_range_loop, reason = "dimension clarity")] for x in 0..Tile::WIDTH as usize { @@ -213,26 +322,26 @@ // Push out the strip if we're moving to a next strip. if !prev_tile.same_loc(&tile) && !prev_tile.prev_loc(&tile) { - debug_assert_eq!( - (prev_tile.x + 1) * Tile::WIDTH - strip.x, - ((alpha_buf.len() - strip.alpha_idx() as usize) / usize::from(Tile::HEIGHT)) as u16, - "The number of columns written to the alpha buffer should equal the number of columns spanned by this strip." - ); - strip_buf.push(strip); + // debug_assert_eq!( + // (prev_tile.x + 1) * Tile::WIDTH - strip.x, + // ((alpha_buf.len() - strip.alpha_idx() as usize) / usize::from(Tile::HEIGHT)) as u16, + // "The number of columns written to the alpha buffer should equal the number of columns spanned by this strip." + // ); + //strip_buf.push(strip); let is_sentinel = tile_idx == tiles.len() as usize; if !prev_tile.same_row(&tile) { // Emit a final strip in the row if there is non-zero winding for the sparse fill, // or unconditionally if we've reached the sentinel tile to end the path (the // `alpha_idx` field is used for width calculations). - if winding_delta != 0 || is_sentinel { - strip_buf.push(Strip::new( - u16::MAX, - prev_tile.y * Tile::HEIGHT, - alpha_buf.len() as u32, - should_fill(winding_delta), - )); - } + // if winding_delta != 0 || is_sentinel { + // strip_buf.push(Strip::new( + // u16::MAX, + // prev_tile.y * Tile::HEIGHT, + // alpha_buf.len() as u32, + // should_fill(winding_delta), + // )); + // } winding_delta = 0; accumulated_winding = f32x4::splat(s, 0.0); @@ -405,10 +514,739 @@ // The trapezoidal area enclosed between the line and the right edge of the pixel // square. let area = 0.5 * h * (2. * px_right_x - line_px_right_yx - line_px_left_yx); + let local_contribution = acc.madd(sign, area); + if x_idx == 0 { + //println!("{}: {:<10.4} ", tile_idx, local_contribution.val[0]); + } + location_winding[x_idx as usize] = location_winding[x_idx as usize] + local_contribution; + acc = acc.madd(sign, h); location_winding[x_idx as usize] += area.madd(sign, acc); acc = h.madd(sign, acc); } - accumulated_winding += acc; + accumulated_winding = accumulated_winding + acc; + } +} + +/// Prepares gpu inputs +fn prepare_gpu_inputs( + tiles: &Tiles, + strip_buf: &mut Vec<Strip>, + pre_merge_buf: &mut Vec<PreMergeTile>, + alpha_buf: &mut Vec<u8>, + fill_rule: Fill, + lines: &[Line], +) { + if tiles.is_empty() { + return; + } + + let should_fill = |winding: i32| match fill_rule { + Fill::NonZero => winding != 0, + Fill::EvenOdd => winding % 2 != 0, + }; + + let mut winding_delta: i32 = 0; + let mut prev_tile = *tiles.get(0); + let mut alpha_offset: u32 = 0; + let initial_alpha_len = alpha_buf.len() as u32; + + const SENTINEL: Tile = Tile::new(u16::MAX, u16::MAX, 0, false); + + let mut start_tile_idx = 0; + let mut start_seg_idx = 0; + let mut strip = Strip::new( + prev_tile.x * Tile::WIDTH, + prev_tile.y * Tile::HEIGHT, + initial_alpha_len, + false, + ); + for (tile_idx, tile) in tiles.iter().copied().chain([SENTINEL]).enumerate() { + let is_start_tile = tile_idx == 0 || !prev_tile.same_loc(&tile); + let is_start_segment = tile_idx == 0 || (is_start_tile && !prev_tile.prev_loc(&tile)); + if is_start_tile { + start_tile_idx = pre_merge_buf.len(); + if tile_idx > 0 { + alpha_offset += (Tile::WIDTH * Tile::HEIGHT) as u32; + } + } + + if is_start_segment { + start_seg_idx = pre_merge_buf.len(); + strip_buf.push(strip); + + let is_sentinel = tile_idx == tiles.len() as usize; + if !prev_tile.same_row(&tile) { + if winding_delta != 0 || is_sentinel { + strip_buf.push(Strip::new( + u16::MAX, + prev_tile.y * Tile::HEIGHT, + initial_alpha_len + alpha_offset, + should_fill(winding_delta), + )); + } + winding_delta = 0; + } + + if is_sentinel { + break; + } + + strip = Strip::new( + tile.x * Tile::WIDTH, + tile.y * Tile::HEIGHT, + initial_alpha_len + alpha_offset, + should_fill(winding_delta), + ); + } + + let line = lines[tile.line_idx() as usize]; + let tile_left_x = f32::from(tile.x) * f32::from(Tile::WIDTH); + let tile_top_y = f32::from(tile.y) * f32::from(Tile::HEIGHT); + let p0_x = line.p0.x - tile_left_x; + let p0_y = line.p0.y - tile_top_y; + let p1_x = line.p1.x - tile_left_x; + let p1_y = line.p1.y - tile_top_y; + + let mut pmt = PreMergeTile { + alpha_index: initial_alpha_len + alpha_offset, + packed_info: 0, + scanned_winding: winding_delta, + padding: 0, + p0: [p0_x, p0_y], + p1: [p1_x, p1_y], + }; + + let sign = (p0_y - p1_y).signum(); + let signed_winding = sign as i32 * tile.winding() as i32; + winding_delta += signed_winding; + + pmt.set_is_end_tile( + tile_idx == tiles.len() as usize - 1 || + !tiles.get(tile_idx as u32 + 1).same_loc(&tile), + ); + pmt.set_is_tile_first_col(tile.x == 0); + + let block_idx = pre_merge_buf.len() >> BLOCK_DIM_LG << BLOCK_DIM_LG; + let start_tile_pack = if start_tile_idx < block_idx { + INVALID_ID + } else { + (start_tile_idx - block_idx) as u32 + }; + pmt.set_tile_start_id(start_tile_pack); + + let start_seg_pack = if start_seg_idx < block_idx { + INVALID_ID + } else { + (start_seg_idx - block_idx) as u32 + }; + pmt.set_seg_start_id(start_seg_pack); + pmt.set_fill_rule(fill_rule == Fill::NonZero); + + pre_merge_buf.push(pmt); + prev_tile = tile; + } + + //TODO alpha buff should be removed in place of a single accumulated u32 + alpha_buf.resize(alpha_buf.len() + alpha_offset as usize, 0); +} + + +#[derive(Debug, Clone, Copy, Default)] +pub struct StitchIndicator { + /// True if the tile's location dependency must be stitched. + pub loc_stitch_required: bool, + /// True if the tile's accumulation dependency must be stitched. + pub acc_stitch_required: bool, + /// The ID of the partition (workgroup) that processed this tile. + pub partition_id: usize, +} + +#[derive(Debug, Clone, Copy, Default)] +pub struct PartitionIndicator { + /// Does this partition contain a `seg_start`? + pub has_seg_start: bool, + /// Does this partition contain a `tile_start`? + pub has_tile_start: bool, +} + +/// Simulate the merge_shader using the PreMergeTile intermediate representation. +/// Note: THIS WILL NOT WORK with the current prepare_gpu_inputs because it is expecting local +/// indexes relative to the call of strip::render, rather than the current global indexes provided +/// by using the pre_merge_buffer length. +fn cpu_merge( + pmt_buf: &mut Vec<PreMergeTile>, + alpha_buf: &mut Vec<u8>, + fill_rule: Fill, + winding_fine_comp: &mut Vec<[[f32; 4]; 4]>, + winding_acc_comp: &mut Vec<[f32; 4]>, + winding_coarse_comp: &mut Vec<i32>, +) { + if pmt_buf.is_empty() { + return; + } + + let pmt_count = pmt_buf.len(); + let mut temp_acc = vec![[0f32; Tile::HEIGHT as usize]; pmt_count]; + let mut temp_fine = vec![[[0f32; Tile::HEIGHT as usize]; Tile::WIDTH as usize]; pmt_count]; + + for gid in 0..pmt_buf.len() { + let pmt = pmt_buf[gid]; + let p0_x = pmt.p0[0]; + let p0_y = pmt.p0[1]; + let p1_x = pmt.p1[0]; + let p1_y = pmt.p1[1]; + + if (p0_y - p1_y).abs() < 1e-6 { + continue; + } + + let (line_top_y, line_top_x, line_bottom_y, line_bottom_x) = if p0_y < p1_y { + (p0_y, p0_x, p1_y, p1_x) + } else { + (p1_y, p1_x, p0_y, p0_x) + }; + let (line_left_x, line_left_y, line_right_x) = if p0_x < p1_x { + (p0_x, p0_y, p1_x) + } else { + (p1_x, p1_y, p0_x) + }; + + + let dx = line_bottom_x - line_top_x; // this becomes the operative issue. + let is_vertical = dx.abs() < 1e-6; // Now use an epsilon to be sure + let dy = line_bottom_y - line_top_y; // we skip horizontal or close to horizontal lines by now + let y_slope = if dx == 0.0 { f32::MAX } else { dy / dx }; + let x_slope = if dy == 0.0 { f32::MAX } else { dx / dy }; + let sign = (p0_y - p1_y).signum(); + + if pmt.is_tile_first_col() && line_left_x < 0. { + let (ymin, ymax) = if is_vertical { + (line_top_y, line_bottom_y) + } else { + let line_viewport_left_y = (line_top_y - line_top_x * y_slope) + .max(line_top_y) + .min(line_bottom_y); + ( + f32::min(line_left_y, line_viewport_left_y), + f32::max(line_left_y, line_viewport_left_y), + ) + }; + + let px_top_y_arr = [0.0, 1.0, 2.0, 3.0]; + let mut h = [0.0; Tile::HEIGHT as usize]; + for y in 0..Tile::HEIGHT as usize { + let px_top_y: f32 = px_top_y_arr[y]; + let px_bottom_y: f32 = 1.0 + px_top_y; + let ymin_clamped = px_top_y.max(ymin); + let ymax_clamped = px_bottom_y.min(ymax); + h[y] = (ymax_clamped - ymin_clamped).max(0.0); + temp_acc[gid][y] = sign * h[y]; + } + + if line_right_x < 0. { + for x_idx in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + temp_fine[gid][x_idx][y] = temp_acc[gid][y]; + } + } + continue; + } + } + + if is_vertical { + let line_x = line_top_x; + let event_x_idx = line_x.floor() as usize; + for x_idx in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + let mut h = 0.0; + let mut area = 0.0; + if x_idx == event_x_idx { + let px_top_y = y as f32; + let px_bottom_y = 1.0 + px_top_y; + + let ymin = line_top_y.max(px_top_y); + let ymax = line_bottom_y.min(px_bottom_y); + let coverage_right = (x_idx as f32 + 1.0) - line_x; + h = (ymax - ymin).max(0.0); + area = h * coverage_right; + } + temp_fine[gid][x_idx][y] = temp_acc[gid][y] + sign * area; + temp_acc[gid][y] += sign * h; + } + } + } else { + for x_idx in 0..Tile::WIDTH { + let x_idx_f = x_idx as f32; + let px_left_x = x_idx_f; + let px_right_x = 1.0 + x_idx_f; + + for y in 0..Tile::HEIGHT as usize { + let px_top_y = y as f32; + let px_bottom_y = 1.0 + px_top_y; + let ymin = line_top_y.max(px_top_y); + let ymax = line_bottom_y.min(px_bottom_y); + + let line_px_left_y = (line_top_y + (px_left_x - line_top_x) * y_slope) + .max(ymin) + .min(ymax); + let line_px_right_y = (line_top_y + (px_right_x - line_top_x) * y_slope) + .max(ymin) + .min(ymax); + + let line_px_left_yx = line_top_x + (line_px_left_y - line_top_y) * x_slope; + let line_px_right_yx = line_top_x + (line_px_right_y - line_top_y) * x_slope; + let h = (line_px_right_y - line_px_left_y).abs(); + let area = 0.5 * h * (2. * px_right_x - line_px_right_yx - line_px_left_yx); + + temp_fine[gid][x_idx as usize][y] = temp_acc[gid][y] + sign * area; + temp_acc[gid][y] += sign * h; + } + } + } + } + + const BLOCK_DIM: usize = 256; + let GRID_DIM: usize = (pmt_buf.len() + BLOCK_DIM - 1) / BLOCK_DIM; + let mut stitch_indicator: Vec<StitchIndicator> = Vec::new(); + let mut part_indicator: Vec<PartitionIndicator> = Vec::new(); + let mut stitch_loc: Vec<[[f32; Tile::HEIGHT as usize]; Tile::WIDTH as usize]> = Vec::new(); + let mut part_loc: Vec<[[f32; Tile::HEIGHT as usize]; Tile::WIDTH as usize]> = Vec::new(); + let mut part_acc_seg: Vec<[f32; Tile::HEIGHT as usize]> = Vec::new(); + let mut part_acc_mine: Vec<[f32; Tile::HEIGHT as usize]> = Vec::new(); + + for wgid in 0..GRID_DIM { + let mut wg_loc = [[[0f32; Tile::HEIGHT as usize]; Tile::WIDTH as usize]; BLOCK_DIM]; + let mut wg_acc = [[0f32; Tile::HEIGHT as usize]; BLOCK_DIM]; + + // dirty, blend the scan with the stitch preparation + for tid in 0..BLOCK_DIM { + let gid = tid + wgid * BLOCK_DIM; + if gid < pmt_count { + wg_acc[tid] = temp_acc[gid]; + wg_loc[tid] = temp_fine[gid]; + } + + if tid > 0 { + for y in 0..Tile::HEIGHT as usize { + wg_acc[tid][y] += wg_acc[tid - 1][y]; + } + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + wg_loc[tid][x][y] += wg_loc[tid - 1][x][y]; + } + } + } + // barrier + + // now duplicate previous logic. . . + if gid < pmt_count { + let pmt = pmt_buf[gid]; + if pmt.is_end_tile() || tid == BLOCK_DIM - 1 { + let mut my_acc = wg_acc[tid]; + let mut my_loc = wg_loc[tid]; + + let tile_start_id = pmt.tile_start_id() as usize; + let seg_start_id = pmt.seg_start_id() as usize; + let tile_start_valid = tile_start_id != INVALID_ID as usize; + let seg_start_valid = seg_start_id != INVALID_ID as usize; + + if tile_start_valid { + if tile_start_id != 0 { + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + my_loc[x][y] -= wg_loc[tile_start_id - 1][x][y]; + } + } + + my_acc = wg_acc[tile_start_id - 1]; + } else { + my_acc = [0f32; 4]; + } + } + + if seg_start_valid { + let scanned_winding = pmt_buf[wgid * BLOCK_DIM + seg_start_id].scanned_winding as f32; + if seg_start_id != 0 { + for y in 0..Tile::HEIGHT as usize { + my_acc[y] += scanned_winding - wg_acc[seg_start_id - 1][y]; + } + } else { + for y in 0..Tile::HEIGHT as usize { + my_acc[y] += scanned_winding; + } + } + } + + if tile_start_valid { + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + my_loc[x][y] += my_acc[y]; + } + } + } + + let mut stitch = StitchIndicator { + loc_stitch_required: false, + acc_stitch_required: false, + partition_id: 0, + }; + + // only end tiles participate in the write out + if pmt.is_end_tile() { + if tile_start_valid { + if seg_start_valid { + temp_fine[gid] = my_loc; + } else { + stitch.acc_stitch_required = true; + stitch.partition_id = wgid; + } + } else { + stitch.loc_stitch_required = true; + stitch.partition_id = wgid; + } + + } + stitch_loc.push(my_loc); + stitch_indicator.push(stitch); + + if tid == BLOCK_DIM - 1 { + // Full OR tile - seg_start + part_acc_mine.push(my_acc); + + // Push Full + let mut seg_acc = [0f32; Tile::HEIGHT as usize]; + if seg_start_valid { + let scanned_winding = pmt_buf[wgid * BLOCK_DIM + seg_start_id].scanned_winding as f32; + if seg_start_id != 0 { + for y in 0..Tile::HEIGHT as usize { + seg_acc[y] = scanned_winding + wg_acc[tid][y] - + wg_acc[seg_start_id - 1][y]; + } + } else { + for y in 0..Tile::HEIGHT as usize { + seg_acc[y] = scanned_winding + wg_acc[tid][y]; + } + } + part_acc_seg.push(seg_acc); + } else { + part_acc_seg.push(wg_acc[tid]); + } + + if tile_start_valid && tile_start_id != 0 { + let mut p_loc = [[0f32; Tile::HEIGHT as usize]; Tile::WIDTH as usize]; + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + p_loc[x][y] += wg_loc[tid][x][y] - wg_loc[pmt.tile_start_id() as usize - 1][x][y]; + } + } + part_loc.push(p_loc); + } else { + part_loc.push(wg_loc[tid]); + } + let p_ind = PartitionIndicator { + has_seg_start: seg_start_valid, + has_tile_start: tile_start_valid, + }; + part_indicator.push(p_ind); + } + } else { + let stitch = StitchIndicator { + loc_stitch_required: false, + acc_stitch_required: false, + partition_id: 0, + }; + stitch_loc.push([[0f32; 4]; 4]); + stitch_indicator.push(stitch); + } + } + } + } + + //Stitching is not done in grid blocked + for gid in 0..pmt_count { + let s_ind = stitch_indicator[gid]; + if s_ind.loc_stitch_required && !s_ind.acc_stitch_required { + let mut loc = stitch_loc[gid]; + let mut lookback_id = s_ind.partition_id - 1; + let mut part_ind:PartitionIndicator; + loop { + // Sum the prev + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + loc[x][y] += part_loc[lookback_id][x][y]; + } + } + + // Did we hit a tile_start? + part_ind = part_indicator[lookback_id]; + if part_ind.has_tile_start { + break; + } else { + lookback_id -= 1; + } + } + + //Once we hit a tile start. . . Then we have to stitch the acc + //Since this partition is guaranteed to have the tile start + let mut acc = part_acc_mine[lookback_id]; + + // If this partition also had a seg start, we're done, otherwise we need to continue + // the traversal + if !part_ind.has_seg_start { + lookback_id -= 1; + loop { + // Sum the prev + for y in 0..Tile::HEIGHT as usize { + acc[y] += part_acc_seg[lookback_id][y]; + } + + // Did we hit a seg_start? + part_ind = part_indicator[lookback_id]; + if part_ind.has_seg_start { + break; + } else { + lookback_id -= 1; + } + } + } + + //Push out + temp_acc[gid] = acc; + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + temp_fine[gid][x][y] = loc[x][y] + acc[y]; + } + } + } + + if s_ind.acc_stitch_required && !s_ind.loc_stitch_required { + let mut acc = [0f32; Tile::HEIGHT as usize]; + let mut lookback_id = s_ind.partition_id - 1; + loop { + let s_acc = part_acc_seg[lookback_id]; + for y in 0..Tile::HEIGHT as usize { + acc[y] += s_acc[y]; + } + if part_indicator[lookback_id].has_seg_start { + break; + } else { + lookback_id -= 1; + } + } + + // Push out + let loc = stitch_loc[gid]; + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + temp_fine[gid][x][y] = loc[x][y] + acc[y]; + } + } + } + } + + for gid in 0..(pmt_count - 1) { + let pmt = pmt_buf[gid]; + if pmt.is_end_tile() { + winding_fine_comp.push(temp_fine[gid]); + winding_acc_comp.push(temp_acc[gid + 1]); + winding_coarse_comp.push(pmt_buf[gid + 1].scanned_winding); + + match fill_rule { + Fill::NonZero => { + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + let area = temp_fine[gid][x][y]; + let coverage = area.abs(); + let mulled = 0.5 + coverage * 255.0; + temp_fine[gid][x][y] = mulled.min(255.0); + } + } + } + Fill::EvenOdd => { + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + let area = temp_fine[gid][x][y]; + let im1 = (0.5 + area * 0.5).floor(); + let coverage = (im1 + area * -2.0).abs(); + let mulled = 0.5 + coverage * 255.0; + temp_fine[gid][x][y] = mulled.min(255.0); + } + } + } + }; + + let mut u8_vals = [0u8; 16]; + let mut i = 0; + for x in 0..Tile::WIDTH as usize { + for y in 0..Tile::HEIGHT as usize { + u8_vals[i] = temp_fine[gid][x][y].round() as u8; + i += 1; + } + } + + // this will be indexed into + alpha_buf.extend_from_slice(&u8_vals); + } + } +} + +//////////////////////////////////////////////////////////////////////////////////////////////////// +pub fn compare_windings( + winding_fine_ref: &[[[f32; 4]; 4]], + winding_coarse_ref: &[i32], + winding_acc_ref: &[[f32; 4]], + winding_fine_comp: &[[[f32; 4]; 4]], + winding_coarse_comp: &[i32], + winding_acc_comp: &[[f32; 4]], +) -> bool { + let coarse_match = compare_coarse_windings(winding_coarse_ref, winding_coarse_comp); + if !coarse_match { + println!("\nComparison halted due to coarse winding mismatch."); + return false; + } + let fine_match = compare_fine_windings(winding_fine_ref, winding_fine_comp); + if !fine_match { + println!("\nComparison failed at fine winding stage."); + return false; + } + let acc_match = compare_acc_windings(winding_acc_ref, winding_acc_comp); + + acc_match && coarse_match && fine_match +} + +fn compare_coarse_windings(reference: &[i32], comp: &[i32]) -> bool { + if reference.len() != comp.len() { + println!( + "❌ FATAL ERROR: Coarse winding vectors have different lengths! Reference: {}, Comp: {}", + reference.len(), + comp.len() + ); + return false; + } + let mut mismatches = 0; + for (i, (ref_val, comp_val)) in reference.iter().zip(comp.iter()).enumerate() { + if ref_val != comp_val { + if mismatches == 0 { + println!("\n--- Coarse Winding Mismatches Found ---"); + } + println!( + "Mismatch at tile index {}: Reference = {}, Comp = {}", + i, ref_val, comp_val + ); + mismatches += 1; + } + } + if mismatches == 0 { + true + } else { + println!( + "❌ Coarse winding comparison FAILED. Found {} mismatches out of {} tiles.", + mismatches, + reference.len() + ); + false + } +} + +fn compare_fine_windings(reference: &[[[f32; 4]; 4]], comp: &[[[f32; 4]; 4]]) -> bool { + if reference.len() != comp.len() { + println!( + "❌ FATAL ERROR: Fine winding vectors have different lengths! Reference: {}, Comp: {}", + reference.len(), + comp.len() + ); + return false; + } + + let mut total_mismatches = 0; + let mut first_mismatch_tile_index = None; + + for i in 0..reference.len() { + let ref_tile = reference[i]; + let comp_tile = comp[i]; + let mut tile_has_mismatch = false; + for x in 0..4 { + for y in 0..4 { + let ref_val = ref_tile[x][y]; + let comp_val = comp_tile[x][y]; + if (ref_val - comp_val).abs() > 0.01 { + total_mismatches += 1; + tile_has_mismatch = true; + } + } + } + if tile_has_mismatch && first_mismatch_tile_index.is_none() { + first_mismatch_tile_index = Some(i); + } + } + + if total_mismatches == 0 { + true + } else { + println!( + "❌ Fine winding comparison FAILED. Found {} total mismatches.", + total_mismatches + ); + if let Some(i) = first_mismatch_tile_index { + println!( + "\n--- Detailed Mismatch Report for First Failing Tile (Index {}) ---", + i + ); + println!("Column-major format: pixel (x, y)"); + println!("{:<45} {:<45}", "Reference (cpu_merge_ref)", "Comp"); + println!("{:-<45} {:-<45}", "-", "-"); + + let ref_tile = reference[i]; + let comp_tile = comp[i]; + + for y in 0..4 { + let mut ref_row_str = String::new(); + let mut comp_row_str = String::new(); + for x in 0..4 { + ref_row_str.push_str(&format!("({x},{y})={:<+8.4} ", ref_tile[x][y])); + comp_row_str.push_str(&format!("({x},{y})={:<+8.4} ", comp_tile[x][y])); + } + println!("{} {}", ref_row_str.trim_end(), comp_row_str.trim_end()); + } + } + false + } +} + +fn compare_acc_windings(reference: &[[f32; 4]], comp: &[[f32; 4]]) -> bool { + if reference.len() != comp.len() { + println!( + "❌ FATAL ERROR: Accumulator winding vectors have different lengths! Reference: {}, Comp: {}", + reference.len(), + comp.len() + ); + return false; + } + let mut mismatches = 0; + for i in 0..reference.len() { + let ref_tile_acc = reference[i]; + let comp_tile_acc = comp[i]; + for y in 0..4 { + if (ref_tile_acc[y] - comp_tile_acc[y]).abs() > 0.01 { + if mismatches == 0 { + println!("\n--- Accumulator Winding Mismatches Found ---"); + } + println!( + "Mismatch at tile index {} row {}: Reference = {:.4}, Comp = {:.4}", + i, y, ref_tile_acc[y], comp_tile_acc[y] + ); + mismatches += 1; + } + } + } + if mismatches == 0 { + true + } else { + println!( + "❌ Accumulator winding comparison FAILED. Found {} mismatches.", + mismatches, + ); + false } }
diff --git a/sparse_strips/vello_common/src/strip_generator.rs b/sparse_strips/vello_common/src/strip_generator.rs index e7bae82..2198c98 100644 --- a/sparse_strips/vello_common/src/strip_generator.rs +++ b/sparse_strips/vello_common/src/strip_generator.rs
@@ -7,7 +7,7 @@ use crate::flatten::{FlattenCtx, Line}; use crate::kurbo::{Affine, PathEl, Stroke}; use crate::peniko::Fill; -use crate::strip::Strip; +use crate::strip::{PreMergeTile, Strip}; use crate::tile::Tiles; use crate::{flatten, strip}; use alloc::vec::Vec; @@ -20,6 +20,9 @@ pub strips: Vec<Strip>, /// The alphas in the storage. pub alphas: Vec<u8>, + /// For the pre merge tiles + pub pre_merge_tiles: Vec<PreMergeTile>, + generation_mode: GenerationMode, } @@ -141,6 +144,7 @@ self.level, &self.tiles, &mut strip_storage.strips, + &mut strip_storage.pre_merge_tiles, &mut strip_storage.alphas, fill_rule, aliasing_threshold,
diff --git a/sparse_strips/vello_cpu/src/render.rs b/sparse_strips/vello_cpu/src/render.rs index ebcb5ef..e5e497b 100644 --- a/sparse_strips/vello_cpu/src/render.rs +++ b/sparse_strips/vello_cpu/src/render.rs
@@ -27,7 +27,7 @@ use vello_common::peniko::{BlendMode, Compose, Fill, Mix}; use vello_common::pixmap::Pixmap; use vello_common::recording::{PushLayerCommand, Recordable, Recorder, Recording, RenderCommand}; -use vello_common::strip::Strip; +use vello_common::strip::{PreMergeTile, Strip}; use vello_common::strip_generator::{GenerationMode, StripGenerator, StripStorage}; #[cfg(feature = "text")] use vello_common::{ @@ -691,8 +691,8 @@ } fn execute_recording(&mut self, recording: &Recording) { - let (cached_strips, cached_alphas) = recording.get_cached_strips(); - let adjusted_strips = self.prepare_cached_strips(cached_strips, cached_alphas); + let (cached_strips, cached_pmt, cached_alphas) = recording.get_cached_strips(); + let adjusted_strips = self.prepare_cached_strips(cached_strips, cached_pmt, cached_alphas); // Use pre-calculated strip start indices from when we generated the cache. let strip_start_indices = recording.get_strip_start_indices(); @@ -910,18 +910,23 @@ fn prepare_cached_strips( &mut self, cached_strips: &[Strip], + cached_pmt: &[PreMergeTile], cached_alphas: &[u8], ) -> Vec<Strip> { // Calculate offset for alpha indices based on current dispatcher's alpha buffer size. let alpha_offset = { let storage = self.dispatcher.strip_storage_mut(); let offset = storage.alphas.len() as u32; - // Extend the dispatcher's alpha buffer with cached alphas. - storage.alphas.extend(cached_alphas); + storage.alphas.extend_from_slice(cached_alphas); + storage.pre_merge_tiles.extend(cached_pmt.iter().map(|pmt| { + let mut adjusted_pmt = *pmt; + adjusted_pmt.alpha_index += offset; + adjusted_pmt + })); offset }; - // Create adjusted strips with corrected alpha indices. + cached_strips .iter() .map(move |strip| {
diff --git a/sparse_strips/vello_hybrid/examples/winit/Cargo.toml b/sparse_strips/vello_hybrid/examples/winit/Cargo.toml index 23e8426..20a0dcd 100644 --- a/sparse_strips/vello_hybrid/examples/winit/Cargo.toml +++ b/sparse_strips/vello_hybrid/examples/winit/Cargo.toml
@@ -16,3 +16,5 @@ vello_hybrid = { workspace = true } vello_example_scenes = { workspace = true } pollster = { workspace = true } +futures-intrusive = "0.5" +bytemuck = { version = "1.15", features = ["derive"] } \ No newline at end of file
diff --git a/sparse_strips/vello_hybrid/examples/winit/src/main.rs b/sparse_strips/vello_hybrid/examples/winit/src/main.rs index 5d429ff..8843a8c 100644 --- a/sparse_strips/vello_hybrid/examples/winit/src/main.rs +++ b/sparse_strips/vello_hybrid/examples/winit/src/main.rs
@@ -22,6 +22,8 @@ window::{Window, WindowId}, }; +use futures_intrusive::channel::shared::oneshot_channel; + const ZOOM_STEP: f64 = 0.1; struct App<'s> { @@ -306,9 +308,58 @@ ) .unwrap(); + let renderer = self.renderers[surface.dev_id].as_mut().unwrap(); + let source_buffer = renderer.get_stitch_indicator_buffer(); + let buffer_size = source_buffer.size(); + + let staging_buffer = device_handle.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Stitch Indicator Readback Buffer"), + size: buffer_size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + encoder.copy_buffer_to_buffer( + source_buffer, + 0, // source offset + &staging_buffer, + 0, // destination offset + buffer_size, + ); + device_handle.queue.submit([encoder.finish()]); surface_texture.present(); + let buffer_slice = staging_buffer.slice(..); + let (sender, receiver) = oneshot_channel(); + buffer_slice.map_async(wgpu::MapMode::Read, move |result| { + sender.send(result).unwrap(); + }); + + let _ = device_handle.device.poll(wgpu::PollType::Wait); + + // if let Some(Ok(())) = pollster::block_on(receiver.receive()) { + // let data = buffer_slice.get_mapped_range(); + // let u32_data: &[u32] = bytemuck::cast_slice(&data); + // let num_to_print = u32_data.len().min(256); + // println!("--- Readback Buffer (First {} u32s) ---", num_to_print); + // for (i, val) in u32_data.iter().take(num_to_print).enumerate() { + // println!("{}: {:<10} ", i, val) + // } + // println!("\n--- End of Readback ---"); + // } + + // if let Some(Ok(())) = pollster::block_on(receiver.receive()) { + // let data = buffer_slice.get_mapped_range(); + // let f32_data: &[f32] = bytemuck::cast_slice(&data); + // let num_to_print = f32_data.len().min(128); + // println!("--- Readback Buffer (First {} f32s) ---", num_to_print); + // for (i, val) in f32_data.iter().take(num_to_print).enumerate() { + // println!("{}: {:<10.4} ", i, val); + // } + // println!("\n--- End of Readback ---"); + // } + device_handle.device.poll(wgpu::PollType::Poll).unwrap(); } _ => {}
diff --git a/sparse_strips/vello_hybrid/src/lib.rs b/sparse_strips/vello_hybrid/src/lib.rs index 1865cc4..87c1e27 100644 --- a/sparse_strips/vello_hybrid/src/lib.rs +++ b/sparse_strips/vello_hybrid/src/lib.rs
@@ -29,7 +29,7 @@ //! //! See the individual module documentation for more details on usage and implementation. -#![no_std] +//#![no_std] extern crate alloc;
diff --git a/sparse_strips/vello_hybrid/src/render/common.rs b/sparse_strips/vello_hybrid/src/render/common.rs index 253445d..bbe9aa0 100644 --- a/sparse_strips/vello_hybrid/src/render/common.rs +++ b/sparse_strips/vello_hybrid/src/render/common.rs
@@ -44,6 +44,20 @@ pub alphas_tex_width_bits: u32, } +/// Configuration for the Merge and Stitch compute shaders +#[repr(C)] +#[derive(Debug, Copy, Clone, Pod, Zeroable)] +pub struct ComputeConfig { + /// PreMergeTileCount; 1:1 with tiles + pub pmt_count: u32, + /// The total number writing tiles. alpha_buffer.len() / (tile::WIDTH * tile::HEIGHT) + pub end_tile_count: u32, + /// Placeholder + pub c: u32, + /// Placeholder + pub d: u32, +} + /// Represents a GPU strip for rendering. /// /// This struct corresponds to the `StripInstance` struct in the shader.
diff --git a/sparse_strips/vello_hybrid/src/render/wgpu.rs b/sparse_strips/vello_hybrid/src/render/wgpu.rs index a862776..51756e4 100644 --- a/sparse_strips/vello_hybrid/src/render/wgpu.rs +++ b/sparse_strips/vello_hybrid/src/render/wgpu.rs
@@ -22,6 +22,7 @@ use alloc::{sync::Arc, vec}; use core::{fmt::Debug, mem, num::NonZeroU64}; use wgpu::Extent3d; +use std::println; use crate::AtlasConfig; use crate::multi_atlas::AtlasId; @@ -32,7 +33,7 @@ render::{ Config, common::{ - GPU_ENCODED_IMAGE_SIZE_TEXELS, GPU_LINEAR_GRADIENT_SIZE_TEXELS, + ComputeConfig, GPU_ENCODED_IMAGE_SIZE_TEXELS, GPU_LINEAR_GRADIENT_SIZE_TEXELS, GPU_RADIAL_GRADIENT_SIZE_TEXELS, GPU_SWEEP_GRADIENT_SIZE_TEXELS, GpuEncodedImage, GpuEncodedPaint, GpuLinearGradient, GpuRadialGradient, GpuSweepGradient, pack_image_offset, pack_image_params, pack_image_size, pack_radial_kind_and_swapped, @@ -50,6 +51,7 @@ paint::ImageSource, peniko, pixmap::Pixmap, + strip::PreMergeTile, tile::Tile, }; use wgpu::{ @@ -99,6 +101,11 @@ Self::new_with(device, render_target_config, RenderSettings::default()) } + /// TODO + pub fn get_stitch_indicator_buffer(&self) -> &wgpu::Buffer { + &self.programs.resources.stitch_indicator_buffer + } + /// Creates a new renderer with specific settings. pub fn new_with( device: &Device, @@ -146,12 +153,41 @@ self.programs.prepare( device, queue, + &scene.strip_storage.pre_merge_tiles, &mut self.gradient_cache, &self.encoded_paints, &scene.strip_storage.alphas, render_size, &self.paint_idxs, ); + + const BLOCK_DIM: u32 = 256; + { + let tile_count = scene.strip_storage.pre_merge_tiles.len() as u32; + println!("PMT Size {}\n", tile_count); + println!("Alpha Size {}\n", scene.strip_storage.alphas.len()); + let workgroup_x = (tile_count + BLOCK_DIM - 1) / BLOCK_DIM; + + let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("Merge Compute Pass"), + timestamp_writes: None, + }); + compute_pass.set_pipeline(&self.programs.merge_pipeline); + compute_pass.set_bind_group(0, &self.programs.resources.merge_bind_group, &[]); + compute_pass.dispatch_workgroups(workgroup_x, 1, 1); + } + { + let end_tile_count = (scene.strip_storage.alphas.len() / 16) as u32; + let workgroup_x = (end_tile_count + BLOCK_DIM - 1) / BLOCK_DIM; + let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("Stitch Compute Pass"), + timestamp_writes: None, + }); + compute_pass.set_pipeline(&self.programs.stitch_pipeline); + compute_pass.set_bind_group(0, &self.programs.resources.merge_bind_group, &[]); + compute_pass.dispatch_workgroups(workgroup_x, 1, 1); + } + let mut junk = RendererContext { programs: &mut self.programs, device, @@ -447,10 +483,19 @@ gradient_bind_group_layout: BindGroupLayout, /// Bind group layout for atlas textures atlas_bind_group_layout: BindGroupLayout, + /// Bind group layout for merge and stitch shaders + merge_bind_group_layout: BindGroupLayout, + /// Pipeline for clearing slots in slot textures. clear_pipeline: RenderPipeline, /// Pipeline for clearing atlas regions. atlas_clear_pipeline: RenderPipeline, + + /// Merge compute pipeline + merge_pipeline: wgpu::ComputePipeline, + /// Stitch compute pipeline + stitch_pipeline: wgpu::ComputePipeline, + /// GPU resources for rendering (created during prepare) resources: GpuResources, /// Dimensions of the rendering target @@ -487,6 +532,16 @@ view_config_buffer: Buffer, /// Config buffer for rendering wide tile commands into a slot texture. slot_config_buffer: Buffer, + /// Config buffer for Stitch and Merge compute shaders. + compute_config_buffer: Buffer, + + // Merge and stitch + premerge_tile_buffer: Buffer, + stitch_indicator_buffer: Buffer, + stitch_loc_buffer: Buffer, + part_indicator_buffer: Buffer, + part_acc_buffer: Buffer, + part_loc_buffer: Buffer, /// Buffer for slot indices used in `clear_slots` clear_slot_indices_buffer: Buffer, @@ -497,6 +552,8 @@ /// Bind group for clear slots operation clear_bind_group: BindGroup, + + merge_bind_group: BindGroup, } const SIZE_OF_CONFIG: NonZeroU64 = NonZeroU64::new(size_of::<Config>() as u64).unwrap(); @@ -633,6 +690,93 @@ }], }); + let merge_bind_group_layout = + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("Merge/Stitch Bind Group Layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, // config + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, // pmt_in + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, // stitch_indicator + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 3, // part_indicator + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 4, // stitch_loc + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 5, // part_acc + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 6, // part_loc + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 7, // output (alphas texture) + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: wgpu::TextureFormat::Rgba32Uint, + view_dimension: wgpu::TextureViewDimension::D2, + }, + count: None, + }, + ], + }); + let strip_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("Strip Shader"), source: wgpu::ShaderSource::Wgsl(vello_sparse_shaders::wgsl::RENDER_STRIPS.into()), @@ -643,6 +787,9 @@ source: wgpu::ShaderSource::Wgsl(vello_sparse_shaders::wgsl::CLEAR_SLOTS.into()), }); + let merge_shader = + device.create_shader_module(wgpu::include_wgsl!("../shaders/merge.wgsl")); + let strip_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("Strip Pipeline Layout"), @@ -770,6 +917,31 @@ cache: None, }); + let merge_pipeline_layout = + device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("Merge/Stitch Pipeline Layout"), + bind_group_layouts: &[&merge_bind_group_layout], + push_constant_ranges: &[], + }); + + let merge_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("Merge Pipeline"), + layout: Some(&merge_pipeline_layout), + module: &merge_shader, + entry_point: Some("merge"), + compilation_options: Default::default(), + cache: None, + }); + + let stitch_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("Stitch Pipeline"), + layout: Some(&merge_pipeline_layout), + module: &merge_shader, + entry_point: Some("stitch"), + compilation_options: Default::default(), + cache: None, + }); + let slot_texture_views: [TextureView; 2] = core::array::from_fn(|_| { device .create_texture(&wgpu::TextureDescriptor { @@ -896,7 +1068,54 @@ &slot_texture_views, ); - let resources = GpuResources { + let compute_config_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("Compute Config Buffer"), + contents: bytemuck::bytes_of(&ComputeConfig { + pmt_count: 0, + end_tile_count: 0, + c: 0, + d: 0, + }), + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }); + let premerge_tile_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Pre-merge Tile Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let stitch_indicator_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Stitch Indicator Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let stitch_loc_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Stitch Location Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let part_indicator_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Partition Indicator Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let part_acc_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Partition Accumulator Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let part_loc_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Partition Location Buffer"), + size: 128, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + + let mut resources = GpuResources { strips_buffer: Self::create_strips_buffer(device, 0), clear_slot_indices_buffer, slot_texture_views, @@ -912,13 +1131,35 @@ gradient_texture, gradient_bind_group, view_config_buffer, + compute_config_buffer, + premerge_tile_buffer, + stitch_indicator_buffer, + stitch_loc_buffer, + part_indicator_buffer, + part_acc_buffer, + part_loc_buffer, + merge_bind_group: Self::create_placeholder_bind_group(device, &merge_bind_group_layout), }; + resources.merge_bind_group = Self::create_merge_bind_group( + device, + &merge_bind_group_layout, + &resources.compute_config_buffer, + &resources.premerge_tile_buffer, + &resources.stitch_indicator_buffer, + &resources.part_indicator_buffer, + &resources.stitch_loc_buffer, + &resources.part_acc_buffer, + &resources.part_loc_buffer, + &resources.alphas_texture, + ); + Self { strip_pipeline, strip_bind_group_layout, encoded_paints_bind_group_layout, gradient_bind_group_layout, + merge_bind_group_layout, atlas_bind_group_layout, resources, alpha_data, @@ -929,9 +1170,110 @@ }, clear_pipeline, atlas_clear_pipeline, + merge_pipeline, + stitch_pipeline, } } + fn create_placeholder_bind_group( + device: &wgpu::Device, + layout: &wgpu::BindGroupLayout +) -> wgpu::BindGroup { + // Create dummy resources that match the binding types + let dummy_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Dummy Uniform"), + size: 16, + usage: wgpu::BufferUsages::UNIFORM, + mapped_at_creation: false, + }); + let dummy_storage_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("Dummy Storage"), + size: 16, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + let dummy_texture = device.create_texture(&wgpu::TextureDescriptor { + label: Some("Dummy Texture"), + size: wgpu::Extent3d { width: 1, height: 1, depth_or_array_layers: 1 }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Rgba32Uint, + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING, + view_formats: &[], + }); + let dummy_texture_view = dummy_texture.create_view(&Default::default()); + + device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("Placeholder Merge Bind Group"), + layout, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: dummy_uniform_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 4, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 5, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 6, resource: dummy_storage_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::TextureView(&dummy_texture_view) }, + ], + }) +} + + fn create_merge_bind_group( + device: &wgpu::Device, + layout: &wgpu::BindGroupLayout, + config: &wgpu::Buffer, + pmt: &wgpu::Buffer, + stitch_indicator: &wgpu::Buffer, + part_indicator: &wgpu::Buffer, + stitch_loc: &wgpu::Buffer, + part_acc: &wgpu::Buffer, + part_loc: &wgpu::Buffer, + alphas: &wgpu::Texture, + ) -> wgpu::BindGroup { + device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("Merge/Stitch Bind Group"), + layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: config.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: pmt.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: stitch_indicator.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 3, + resource: part_indicator.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 4, + resource: stitch_loc.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 5, + resource: part_acc.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 6, + resource: part_loc.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 7, + resource: wgpu::BindingResource::TextureView( + &alphas.create_view(&Default::default()), + ), + }, + ], + }) + } + fn create_strips_buffer(device: &Device, required_strips_size: u64) -> Buffer { device.create_buffer(&wgpu::BufferDescriptor { label: Some("Strips Buffer"), @@ -979,7 +1321,9 @@ sample_count: 1, dimension: wgpu::TextureDimension::D2, format: wgpu::TextureFormat::Rgba32Uint, - usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST, + usage: wgpu::TextureUsages::TEXTURE_BINDING + | wgpu::TextureUsages::COPY_DST + | wgpu::TextureUsages::STORAGE_BINDING, view_formats: &[], }) } @@ -1172,18 +1516,99 @@ &mut self, device: &Device, queue: &Queue, + pmt: &[PreMergeTile], gradient_cache: &mut GradientRampCache, encoded_paints: &[GpuEncodedPaint], alphas: &[u8], new_render_size: &RenderSize, paint_idxs: &[u32], ) { + let mut new_bg = false; let max_texture_dimension_2d = device.limits().max_texture_dimension_2d; - self.maybe_resize_alphas_tex(device, max_texture_dimension_2d, alphas); + new_bg |= self.maybe_resize_alphas_tex(device, max_texture_dimension_2d, alphas); self.maybe_resize_encoded_paints_tex(device, max_texture_dimension_2d, paint_idxs); self.maybe_update_config_buffer(queue, max_texture_dimension_2d, new_render_size); - self.upload_alpha_texture(queue, alphas); + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.premerge_tile_buffer, + pmt.len() as u64 * std::mem::size_of::<PreMergeTile>() as u64, + "PreMergeTile Buffer", + ); + + let end_tiles = (alphas.len() / 16) as u64; + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.stitch_indicator_buffer, + end_tiles * 4, + "StitchIndicator", + ); + + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.stitch_loc_buffer, + end_tiles * 64, + "StitchLoc", + ); + + const BLOCK_DIM: u64 = 256; + let num_partitions = (pmt.len() as u64 + BLOCK_DIM - 1) / BLOCK_DIM; + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.part_indicator_buffer, + num_partitions * 4, + "PartIndicator", + ); + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.part_acc_buffer, + num_partitions * 32, + "PartAcc", + ); + new_bg |= Self::maybe_resize_buffer( + device, + &mut self.resources.part_loc_buffer, + num_partitions * 64, + "PartLoc", + ); + + if new_bg { + self.resources.merge_bind_group = Self::create_merge_bind_group( + device, + &self.merge_bind_group_layout, + &self.resources.compute_config_buffer, + &self.resources.premerge_tile_buffer, + &self.resources.stitch_indicator_buffer, + &self.resources.part_indicator_buffer, + &self.resources.stitch_loc_buffer, + &self.resources.part_acc_buffer, + &self.resources.part_loc_buffer, + &self.resources.alphas_texture, + ); + } + + //TODO maybe_update_compute_config_buffer + queue.write_buffer( + &self.resources.compute_config_buffer, + 0, + bytemuck::bytes_of(&ComputeConfig { + pmt_count: pmt.len() as u32, + end_tile_count: alphas.len() as u32 / 16, + c: 0, + d: 0, + }), + ); + + if !pmt.is_empty() { + queue.write_buffer( + &self.resources.premerge_tile_buffer, + 0, + bytemuck::cast_slice(pmt), + ); + } + + + //self.upload_alpha_texture(queue, alphas); //delete this when alpha_buff is removed self.upload_encoded_paints_texture(queue, encoded_paints); if gradient_cache.has_changed() { @@ -1199,7 +1624,7 @@ device: &Device, max_texture_dimension_2d: u32, alphas: &[u8], - ) { + ) -> bool { let required_alpha_height = u32::try_from(alphas.len()) .unwrap() // There are 16 1-byte alpha values per texel. @@ -1239,7 +1664,9 @@ &self.resources.view_config_buffer, &self.resources.slot_texture_views, ); + return true } + false } /// Update the encoded paints texture size if needed. @@ -1345,6 +1772,26 @@ } } + fn maybe_resize_buffer( + device: &Device, + buffer: &mut Buffer, + required_size: u64, + label: &'static str, + ) -> bool { + if required_size > buffer.size() { + let new_size = (required_size + 15) & !15; + *buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some(label), + size: new_size, + usage: buffer.usage(), + mapped_at_creation: false, + }); + true + } else { + false + } + } + /// Resize the texture array to accommodate more atlases. fn maybe_resize_atlas_texture_array( device: &Device,
diff --git a/sparse_strips/vello_hybrid/src/scene.rs b/sparse_strips/vello_hybrid/src/scene.rs index 40234d8..986aced 100644 --- a/sparse_strips/vello_hybrid/src/scene.rs +++ b/sparse_strips/vello_hybrid/src/scene.rs
@@ -16,7 +16,7 @@ use vello_common::peniko::color::palette::css::BLACK; use vello_common::peniko::{BlendMode, Compose, Fill, Mix}; use vello_common::recording::{PushLayerCommand, Recordable, Recorder, Recording, RenderCommand}; -use vello_common::strip::Strip; +use vello_common::strip::{PreMergeTile, Strip}; use vello_common::strip_generator::{GenerationMode, StripGenerator, StripStorage}; use crate::AtlasConfig; @@ -439,8 +439,8 @@ } fn execute_recording(&mut self, recording: &Recording) { - let (cached_strips, cached_alphas) = recording.get_cached_strips(); - let adjusted_strips = self.prepare_cached_strips(cached_strips, cached_alphas); + let (cached_strips, cached_pmt, cached_alphas) = recording.get_cached_strips(); + let adjusted_strips = self.prepare_cached_strips(cached_strips, cached_pmt, cached_alphas); // Use pre-calculated strip start indices from when we generated the cache let strip_start_indices = recording.get_strip_start_indices(); @@ -637,13 +637,23 @@ fn prepare_cached_strips( &mut self, cached_strips: &[Strip], + cached_pmt: &[PreMergeTile], cached_alphas: &[u8], ) -> Vec<Strip> { // Calculate offset for alpha indices based on current buffer size. let alpha_offset = self.strip_storage.alphas.len() as u32; - // Extend current alpha buffer with cached alphas. - self.strip_storage.alphas.extend(cached_alphas); - // Create adjusted strips with corrected alpha indices + + // Extend current alpha and pre-merge tile buffers with cached data. + self.strip_storage.alphas.extend_from_slice(cached_alphas); + self.strip_storage + .pre_merge_tiles + .extend(cached_pmt.iter().map(move |pmt| { + let mut adjusted_pmt = *pmt; + adjusted_pmt.alpha_index += alpha_offset; + adjusted_pmt + })); + + // Create adjusted strips with corrected alpha indices. cached_strips .iter() .map(move |strip| {
diff --git a/sparse_strips/vello_hybrid/src/shaders/merge.wgsl b/sparse_strips/vello_hybrid/src/shaders/merge.wgsl new file mode 100644 index 0000000..a027467 --- /dev/null +++ b/sparse_strips/vello_hybrid/src/shaders/merge.wgsl
@@ -0,0 +1,536 @@ +// Copyright 2025 the Vello Authors +// SPDX-License-Identifier: Apache-2.0 OR MIT +struct Config { + // Count of pre merge tiles + pmt_count: u32, + // Count of tiles which will write to the alpha texture + end_tile_count: u32, + c: u32, + d: u32, +}; + +struct PreMergeTile { + // The index into the alpha buffer that an "end tile" should write to. + alpha_index: u32, + // Contains tile location information + packed_info: u32, + // The exclusive prefix sum of the signed winding number. + // Uploading it saves doing the scan here + scanned_winding: i32, + padding: u32, + // Line points, adjusted by the tile position + p0: vec2f, + p1: vec2f, +}; + +const TILE_HEIGHT = 4u; +const TILE_WIDTH = 4u; + +@group(0) @binding(0) +var<uniform> config : Config; + +@group(0) @binding(1) +var<storage, read> pmt_in: array<PreMergeTile>; + +@group(0) @binding(2) +var<storage, read_write> stitch_indicator: array<u32>; + +@group(0) @binding(3) +var<storage, read_write> part_indicator: array<u32>; + +@group(0) @binding(4) +var<storage, read_write> stitch_loc: array<array<array<f32, TILE_HEIGHT>, TILE_WIDTH>>; + +// This is double strided. This is because we need either: +// 1) The reduction from the last segment start to the end of the partition +// 2) The reduction from the last segment start to the last tile start +// These are not mutually exclusive! Stitching requires both! +@group(0) @binding(5) +var<storage, read_write> part_acc: array<array<f32, TILE_HEIGHT>>; + +@group(0) @binding(6) +var<storage, read_write> part_loc: array<array<array<f32, TILE_HEIGHT>, TILE_WIDTH>>; + +@group(0) @binding(7) +var output: texture_storage_2d<rgba32uint, write>; + +// @group(0) @binding(7) +// var<storage, read_write> alpha_buff: array<vec4<u32>>; + +// MSB LSB +// 31------------------21|20------------------12|11-------------------3|2|1|0| +// | Unused (11) | Seg Start ID (9) | Tile Start ID (9) |F|T|E| +// +// F = FILL_RULE_MASK +// T = IS_TILE_FIRST_COL_MASK +// E = IS_END_TILE_MASK +const IS_END_TILE_MASK: u32 = 1u; +const IS_TILE_FIRST_COL_MASK: u32 = 2u; +const FILL_RULE_MASK: u32 = 4u; +const TILE_START_ID_SHIFT: u32 = 3u; +const SEG_START_ID_SHIFT: u32 = 12u; +const ID_MASK: u32 = 0x1ffu; +const INVALID_ID: u32 = 256u; + +fn is_end_tile(packed_info: u32) -> bool { + return (packed_info & IS_END_TILE_MASK) != 0u; +} + +fn is_tile_first_col(packed_info: u32) -> bool { + return (packed_info & IS_TILE_FIRST_COL_MASK) != 0u; +} + +fn is_fill_rule_non_zero(packed_info: u32) -> bool { + return (packed_info & FILL_RULE_MASK) != 0u; +} + +fn get_tile_start_id(packed_info: u32) -> u32 { + return (packed_info >> TILE_START_ID_SHIFT) & ID_MASK; +} + +fn get_seg_start_id(packed_info: u32) -> u32 { + return (packed_info >> SEG_START_ID_SHIFT) & ID_MASK; +} + +// MSB LSB +// 31|30|29--------------------------------------------------------------0| +// |L|A| Partition ID (30 bits) | +// +// T = STITCH_LOC_MASK +// A = STITCH_ACC_MASK +const STITCH_LOC_MASK = 1u << 31u; +const STITCH_ACC_MASK = 1u << 30u; +const PART_ID_MASK = 0x3fffffffu; + +fn loc_stitch_required(in: u32) -> bool { + return (in & STITCH_LOC_MASK) != 0u; +} + +fn acc_stitch_required(in: u32) -> bool { + return (in & STITCH_ACC_MASK) != 0u; +} + +fn get_part_id(in: u32) -> u32 { + return in & PART_ID_MASK; +} + +// MSB LSB +// 31-----------------------------------------------------------------3|1|0| +// | Unused (30 bits) |T|S| +// +// S = PART_SEG_START_MASK +// T = PART_TILE_START_MASK +const PART_SEG_START_MASK: u32 = 1u << 0u; +const PART_TILE_START_MASK: u32 = 1u << 1u; +fn part_has_seg_start(in: u32) -> bool { + return (in & PART_SEG_START_MASK) != 0u; +} + +fn part_has_tile_start(in: u32) -> bool { + return (in & PART_TILE_START_MASK) != 0u; +} + +const BLOCK_DIM = 256u; +const LG_BLOCK_DIM = 8u; +const SCAN_CONST = LG_BLOCK_DIM - 1u; +const TILE_SIZE = TILE_HEIGHT * TILE_WIDTH; +const EPSILON = 1e-6; + +// TODO! Do not use more than 4096 +var<workgroup> wg_acc: array<array<f32, TILE_HEIGHT>, BLOCK_DIM>; +var<workgroup> wg_loc: array<array<array<f32, TILE_HEIGHT>, TILE_WIDTH>, BLOCK_DIM>; + +// 1 thread per pre_pmt +@compute @workgroup_size(BLOCK_DIM, 1, 1) +fn merge(@builtin(local_invocation_id) tid: vec3<u32>, + @builtin(global_invocation_id) gid: vec3<u32>, + @builtin(workgroup_id) wgid: vec3<u32>) { + var pmt: PreMergeTile; + if (gid.x < config.pmt_count) { + pmt = pmt_in[gid.x]; + } else { + pmt = PreMergeTile( + 0xffffffffu, + 0u, + 0, + 0u, + vec2f(0.0, 0.0), + vec2f(0.0, 0.0), + ); + } + + var acc = array<f32, TILE_HEIGHT>(); + var loc = array<array<f32, TILE_HEIGHT>, TILE_WIDTH>(); + if (abs(pmt.p0.y - pmt.p1.y) >= EPSILON) { // If not horizontal. . . + var line_top_y: f32; + var line_top_x: f32; + var line_bottom_y: f32; + var line_bottom_x: f32; + if (pmt.p0.y < pmt.p1.y) { + line_top_y = pmt.p0.y; + line_top_x = pmt.p0.x; + line_bottom_y = pmt.p1.y; + line_bottom_x = pmt.p1.x; + } else { + line_top_y = pmt.p1.y; + line_top_x = pmt.p1.x; + line_bottom_y = pmt.p0.y; + line_bottom_x = pmt.p0.x; + } + + var line_left_x: f32; + var line_left_y: f32; + var line_right_x: f32; + if (pmt.p0.x < pmt.p1.x) { + line_left_x = pmt.p0.x; + line_left_y = pmt.p0.y; + line_right_x = pmt.p1.x; + } else { + line_left_x = pmt.p1.x; + line_left_y = pmt.p1.y; + line_right_x = pmt.p0.x; + } + + let dx = line_bottom_x - line_top_x; + let dy = line_bottom_y - line_top_y; + let is_vertical = abs(dx) < EPSILON; + let y_slope = dy / dx; + let x_slope = dx / dy; + + // Unnecessary, we remove vertical and horizontal cases, so a comparison will be sufficient; + // i.e. dont care about sign of zero + //let sign = select(1.0f, -1.0f, (bitcast<u32>(pmt.p0.y - pmt.p1.y) & 0x80000000u) != 0u); + let sign = select(1.0f, -1.0f, pmt.p0.y < pmt.p1.y); + + if (is_tile_first_col(pmt.packed_info) && line_left_x < 0.0f) { + var ymin: f32; + var ymax: f32; + if (is_vertical) { + ymin = line_top_y; + ymax = line_bottom_y; + } else { + let line_viewport_left_y = + min(max((line_top_y - line_top_x * y_slope), line_top_y), line_bottom_y); + ymin = min(line_left_y, line_viewport_left_y); + ymax = max(line_left_y, line_viewport_left_y); + } + + var h = array<f32, TILE_HEIGHT>(); + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + let px_top_y = f32(y); + let px_bottom_y = 1.0f + px_top_y; + let ymin_clamped = max(ymin, px_top_y); + let ymax_clamped = min(ymax, px_bottom_y); + h[y] = max(ymax_clamped - ymin_clamped, 0.0f); + acc[y] = sign * h[y]; + } + } else { + // No need to clear in wgsl, but other shader languages required. + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] = 0.0f; + } + } + + if (line_right_x >= 0.0f) { + if (is_vertical) { + let x_int = u32(floor(line_top_x)); + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + var h = 0.0f; + var area = 0.0f; + if x == x_int { + let px_top_y = f32(y); + let px_bottom_y = 1.0f + px_top_y; + let ymin = max(px_top_y, line_top_y); + let ymax = min(px_bottom_y, line_bottom_y); + let coverage_right = f32(x) + 1.0f - line_top_x; + h = max(ymax - ymin, 0.0f); + area = h * coverage_right; + } + loc[x][y] = acc[y] + sign * area; + acc[y] += sign * h; + } + } + } + + if (!is_vertical) { + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + let px_left_x = f32(x); + let px_right_x = 1.0f + px_left_x; + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + let px_top_y = f32(y); + let px_bottom_y = 1.0f + px_top_y; + let ymin = max(line_top_y, px_top_y); + let ymax = min(line_bottom_y, px_bottom_y); + + let line_px_left_y = + min(max((line_top_y + (px_left_x - line_top_x) * y_slope), ymin), ymax); + let line_px_right_y = + min(max((line_top_y + (px_right_x - line_top_x) * y_slope), ymin), ymax); + let line_px_left_yx = + line_top_x + (line_px_left_y - line_top_y) * x_slope; + let line_px_right_yx = + line_top_x + (line_px_right_y - line_top_y) * x_slope; + let h = abs(line_px_right_y - line_px_left_y); + let area = + 0.5f * h * (2.0f * px_right_x - line_px_right_yx - line_px_left_yx); + loc[x][y] = acc[y] + sign * area; + acc[y] += sign * h; + } + } + } + + wg_loc[tid.x] = loc; + } else { + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + wg_loc[tid.x][x] = acc; + } + } + wg_acc[tid.x] = acc; + } else { + // TODO For non wgsl this needs to clear + wg_acc[tid.x] = acc; + wg_loc[tid.x] = loc; + } + + for (var i = 1u; i <= BLOCK_DIM; i <<= 1u) { + workgroupBarrier(); + let ii = tid.x - i; + if (tid.x >= i) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] += wg_acc[ii][y]; + } + + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + loc[x][y] += wg_loc[ii][x][y]; + } + } + } + workgroupBarrier(); + + if (tid.x >= i) { + wg_acc[tid.x] = acc; + wg_loc[tid.x] = loc; + } + } + workgroupBarrier(); + + var scanned_winding = 0.0f; + let end_tile = is_end_tile(pmt.packed_info); + let seg_start_id = get_seg_start_id(pmt.packed_info); + let tile_start_id = get_tile_start_id(pmt.packed_info); + let seg_is_valid = seg_start_id != INVALID_ID; + let tile_is_valid = tile_start_id != INVALID_ID; + if (end_tile || tid.x == BLOCK_DIM - 1u) { + if (tile_is_valid) { + if (tile_start_id != 0u) { + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + loc[x][y] -= wg_loc[tile_start_id - 1u][x][y]; + } + } + acc = wg_acc[tile_start_id - 1u]; + } else { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] = 0.0f; + } + } + } + + if (seg_is_valid) { + scanned_winding = f32(pmt_in[(wgid.x << LG_BLOCK_DIM) + seg_start_id].scanned_winding); + if (seg_start_id != 0u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] += scanned_winding - wg_acc[seg_start_id - 1u][y]; + } + } else { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] += scanned_winding; + } + } + } + + // At this point, if this is the last thread in the threadBlock, this contains the reduction + // from the tile start + if (tid.x == BLOCK_DIM - 1u) { + part_loc[wgid.x] = loc; + } + + // This will add in the acc for the cases we want: tile_is_valid, tile_is_valid && seg_is_valid + if (tile_is_valid) { + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + loc[x][y] += acc[y]; + } + } + } + + // Only end tiles participate in the write out + if (end_tile) { + var s_ind = 0u; + if (tile_is_valid && seg_is_valid) { // Safe to write out + var final_alphas: vec4<u32>; + if (is_fill_rule_non_zero(pmt.packed_info)) { + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + final_alphas[x] = pack4x8unorm(abs(vec4<f32>(loc[x][0], loc[x][1], + loc[x][2], loc[x][3]))); + } + } else { + // EvenOdd fill rule logic + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + let area = vec4<f32>(loc[x][0], loc[x][1], loc[x][2], loc[x][3]); + let im1 = floor(area * 0.5 + 0.5); + let coverage = abs(area - 2.0 * im1); + final_alphas[x] = pack4x8unorm(coverage); + } + } + let tex_dims = textureDimensions(output); + let tex_width = tex_dims.x; + let output_coords = vec2<u32>( + (pmt.alpha_index >> 4) % tex_width, + (pmt.alpha_index >> 4) / tex_width, + ); + if (output_coords.y < tex_dims.y) { + textureStore(output, output_coords, final_alphas); + } + } else { + // TODO fill rule on stitch + s_ind = wgid.x | + select(0u, STITCH_ACC_MASK, tile_is_valid && !seg_is_valid) | + select(0u, STITCH_LOC_MASK, !tile_is_valid); + stitch_loc[pmt.alpha_index >> 4u] = loc; + } + stitch_indicator[pmt.alpha_index >> 4u] = s_ind; + } + } + + if (tid.x == BLOCK_DIM - 1) { + part_acc[wgid.x << 1u] = acc; + var seg_acc = wg_acc[tid.x]; + if (seg_is_valid) { + if (seg_start_id != 0u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + seg_acc[y] += scanned_winding - wg_acc[seg_start_id - 1u][y]; + } + } else { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + seg_acc[y] += scanned_winding; + } + } + } + part_acc[(wgid.x << 1u) + 1u] = seg_acc; + + part_indicator[wgid.x] = select(0u, PART_TILE_START_MASK, tile_is_valid) | + select(0u, PART_SEG_START_MASK, seg_is_valid); + } +} + +// 1 thread per end_tile +// TODO once subgroups, this should be 1 subgroup : tile +@compute @workgroup_size(BLOCK_DIM, 1, 1) +fn stitch(@builtin(local_invocation_id) tid: vec3<u32>, + @builtin(global_invocation_id) gid: vec3<u32>, + @builtin(workgroup_id) wgid: vec3<u32>) { + var s_indicator = 0u; + if (gid.x < config.end_tile_count) { + s_indicator = stitch_indicator[gid.x]; + } + + // No stitching needed or oob + if (s_indicator == 0u) { + return; + } + let part_id = get_part_id(s_indicator); + var loc = stitch_loc[gid.x]; + var acc = array<f32, TILE_HEIGHT>(); + if (loc_stitch_required(s_indicator)) { + var lookback_id = part_id - 1u; + var part_ind: u32; + while (true) { + let p_loc = part_loc[lookback_id]; + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + loc[x][y] += p_loc[x][y]; + } + } + + // Did we hit a tile start? + part_ind = part_indicator[lookback_id]; + if (part_has_tile_start(part_ind)) { + break; + } else { + lookback_id -= 1u; + } + } + + lookback_id <<= 1u; + acc = part_acc[lookback_id]; + + // If the tile start also included a seg start, we're done. Else we will have to traverse to + // the previous seg start + if (!part_has_seg_start(part_ind)) { + lookback_id -= 1u; // Down to the upper + while (true) { + let p_acc = part_acc[lookback_id]; + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] += p_acc[y]; + } + + if (part_has_seg_start(part_indicator[lookback_id >> 1u])) { + break; + } else { + lookback_id -= 2u; + } + } + } + } + + if (acc_stitch_required(s_indicator)) { + var lookback_id = (part_id - 1u) << 1u; + while (true) { + var s_acc = part_acc[lookback_id + 1u]; + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + acc[y] += s_acc[y]; + } + + if (part_has_seg_start(part_indicator[lookback_id >> 1u])) { + break; + } else { + lookback_id -= 2u; + } + } + } + + // Combine acc and loc + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + for (var y = 0u; y < TILE_HEIGHT; y += 1u) { + loc[x][y] += acc[y]; + } + } + + var final_alphas: vec4<u32>; + if (true) { // TODO put fill rule onto the stitch + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + final_alphas[x] = pack4x8unorm(abs(vec4<f32>(loc[x][0], loc[x][1], + loc[x][2], loc[x][3]))); + } + } else { + // EvenOdd fill rule logic + for (var x = 0u; x < TILE_WIDTH; x += 1u) { + let area = vec4<f32>(loc[x][0], loc[x][1], loc[x][2], loc[x][3]); + let im1 = floor(area * 0.5 + 0.5); + let coverage = abs(area - 2.0 * im1); + final_alphas[x] = pack4x8unorm(coverage); + } + } + let tex_dims = textureDimensions(output); + let tex_width = tex_dims.x; + let output_coords = vec2<u32>( + gid.x % tex_width, + gid.x / tex_width, + ); + if (output_coords.y < tex_dims.y) { + textureStore(output, output_coords, final_alphas); + } +}
diff --git a/sparse_strips/vello_toy/src/debug.rs b/sparse_strips/vello_toy/src/debug.rs index 01732a8..f2b390e 100644 --- a/sparse_strips/vello_toy/src/debug.rs +++ b/sparse_strips/vello_toy/src/debug.rs
@@ -33,6 +33,7 @@ let mut line_buf = vec![]; let mut tiles = Tiles::new(Level::new()); let mut strip_buf = vec![]; + let mut pmt_buf = vec![]; let mut alpha_buf = vec![]; let mut wide = Wide::<MODE_CPU>::new(args.width, args.height); @@ -79,6 +80,7 @@ Level::new(), &tiles, &mut strip_buf, + &mut pmt_buf, &mut alpha_buf, args.fill_rule, None,