forked from linebender/vello
-
Notifications
You must be signed in to change notification settings - Fork 0
/
pathtag_scan1.wgsl
66 lines (57 loc) · 1.84 KB
/
pathtag_scan1.wgsl
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense
// This shader computes the scan of reduced tag monoids given
// two levels of reduction.
#import config
#import pathtag
@group(0) @binding(0)
var<storage> reduced: array<TagMonoid>;
@group(0) @binding(1)
var<storage> reduced2: array<TagMonoid>;
@group(0) @binding(2)
var<storage, read_write> tag_monoids: array<TagMonoid>;
let LG_WG_SIZE = 8u;
let WG_SIZE = 256u;
var<workgroup> sh_parent: array<TagMonoid, WG_SIZE>;
// These could be combined?
var<workgroup> sh_monoid: array<TagMonoid, WG_SIZE>;
@compute @workgroup_size(256)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
var agg = tag_monoid_identity();
if local_id.x < wg_id.x {
agg = reduced2[local_id.x];
}
sh_parent[local_id.x] = agg;
for (var i = 0u; i < LG_WG_SIZE; i += 1u) {
workgroupBarrier();
if local_id.x + (1u << i) < WG_SIZE {
let other = sh_parent[local_id.x + (1u << i)];
agg = combine_tag_monoid(agg, other);
}
workgroupBarrier();
sh_parent[local_id.x] = agg;
}
let ix = global_id.x;
agg = reduced[ix];
sh_monoid[local_id.x] = agg;
for (var i = 0u; i < LG_WG_SIZE; i += 1u) {
workgroupBarrier();
if local_id.x >= 1u << i {
let other = sh_monoid[local_id.x - (1u << i)];
agg = combine_tag_monoid(other, agg);
}
workgroupBarrier();
sh_monoid[local_id.x] = agg;
}
workgroupBarrier();
// prefix up to this workgroup
var tm = sh_parent[0];
if local_id.x > 0u {
tm = combine_tag_monoid(tm, sh_monoid[local_id.x - 1u]);
}
// exclusive prefix sum, granularity of 4 tag bytes * workgroup size
tag_monoids[ix] = tm;
}