// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense #import config #import drawtag @group(0) @binding(0) var<uniform> config: Config; @group(0) @binding(1) var<storage> scene: array<u32>; @group(0) @binding(2) var<storage, read_write> reduced: array<DrawMonoid>; let WG_SIZE = 256u; var<workgroup> sh_scratch: array<DrawMonoid, WG_SIZE>; #import util @compute @workgroup_size(256) fn main( @builtin(global_invocation_id) global_id: vec3<u32>, @builtin(local_invocation_id) local_id: vec3<u32>, ) { let ix = global_id.x; let tag_word = read_draw_tag_from_scene(ix); var agg = map_draw_tag(tag_word); sh_scratch[local_id.x] = agg; for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) { workgroupBarrier(); if local_id.x + (1u << i) < WG_SIZE { let other = sh_scratch[local_id.x + (1u << i)]; agg = combine_draw_monoid(agg, other); } workgroupBarrier(); sh_scratch[local_id.x] = agg; } if local_id.x == 0u { reduced[ix >> firstTrailingBit(WG_SIZE)] = agg; } }