// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense #import config #import pathtag @group(0) @binding(0) var config: Config; @group(0) @binding(1) var scene: array; @group(0) @binding(2) var reduced: array; let LG_WG_SIZE = 8u; let WG_SIZE = 256u; var sh_scratch: array; @compute @workgroup_size(256) fn main( @builtin(global_invocation_id) global_id: vec3, @builtin(local_invocation_id) local_id: vec3, ) { let ix = global_id.x; let tag_word = scene[config.pathtag_base + ix]; var agg = reduce_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_tag_monoid(agg, other); } workgroupBarrier(); sh_scratch[local_id.x] = agg; } if local_id.x == 0u { reduced[ix >> LG_WG_SIZE] = agg; } }