blob: 051d8f84ec4174cd4a74ac63687c2010729c789c [file] [log] [blame] [edit]
// 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;
}
}