1import package::realtime::world_cache_query::WORLD_CACHE_EMPTY_CELL;
2import package::realtime::bindings::{
3    world_cache,
4};
5
6enable wgpu_ray_query;
7
8@group(2) @binding(0) var<storage, read_write> world_cache_active_cells_dispatch: vec3<u32>;
9
10var<workgroup> w1: array<u32, 1024u>;
11var<workgroup> w2: array<u32, 1024u>;
12
13@compute @workgroup_size(1024, 1, 1)
14fn decay_world_cache(@builtin(global_invocation_id) global_id: vec3<u32>) {
15    var life = world_cache.life[global_id.x];
16    if life > 0u {
17        life -= 1u;
18        world_cache.life[global_id.x] = life;
19
20        if life == 0u {
21            world_cache.checksums[global_id.x] = WORLD_CACHE_EMPTY_CELL;
22            world_cache.radiance[global_id.x] = vec4(0.0);
23            world_cache.luminance_deltas[global_id.x] = 0.0;
24        }
25    }
26}
27
28@compute @workgroup_size(1024, 1, 1)
29fn compact_world_cache_single_block(
30    @builtin(global_invocation_id) cell_id: vec3<u32>,
31    @builtin(local_invocation_index) t: u32,
32) {
33    if t == 0u { w1[0u] = 0u; } else { w1[t] = u32(world_cache.life[cell_id.x - 1u] != 0u); }; workgroupBarrier();
34    if t < 1u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 1u]; } workgroupBarrier();
35    if t < 2u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 2u]; } workgroupBarrier();
36    if t < 4u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 4u]; } workgroupBarrier();
37    if t < 8u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 8u]; } workgroupBarrier();
38    if t < 16u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 16u]; } workgroupBarrier();
39    if t < 32u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 32u]; } workgroupBarrier();
40    if t < 64u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 64u]; } workgroupBarrier();
41    if t < 128u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 128u]; } workgroupBarrier();
42    if t < 256u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 256u]; } workgroupBarrier();
43    if t < 512u { world_cache.a[cell_id.x] = w2[t]; } else { world_cache.a[cell_id.x] = w2[t] + w2[t - 512u]; }
44}
45
46@compute @workgroup_size(1024, 1, 1)
47fn compact_world_cache_blocks(@builtin(local_invocation_index) t: u32) {
48    if t == 0u { w1[0u] = 0u; } else { w1[t] = world_cache.a[t * 1024u - 1u] + u32(world_cache.life[t * 1024u - 1u] != 0u); }; workgroupBarrier();
49    if t < 1u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 1u]; } workgroupBarrier();
50    if t < 2u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 2u]; } workgroupBarrier();
51    if t < 4u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 4u]; } workgroupBarrier();
52    if t < 8u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 8u]; } workgroupBarrier();
53    if t < 16u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 16u]; } workgroupBarrier();
54    if t < 32u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 32u]; } workgroupBarrier();
55    if t < 64u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 64u]; } workgroupBarrier();
56    if t < 128u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 128u]; } workgroupBarrier();
57    if t < 256u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 256u]; } workgroupBarrier();
58    if t < 512u { world_cache.b[t] = w2[t]; } else { world_cache.b[t] = w2[t] + w2[t - 512u]; }
59}
60
61@compute @workgroup_size(1024, 1, 1)
62fn compact_world_cache_write_active_cells(
63    @builtin(global_invocation_id) cell_id: vec3<u32>,
64    @builtin(workgroup_id) workgroup_id: vec3<u32>,
65    @builtin(local_invocation_index) thread_index: u32,
66) {
67    let compacted_index = world_cache.a[cell_id.x] + world_cache.b[workgroup_id.x];
68    let cell_active = world_cache.life[cell_id.x] != 0u;
69
70    if cell_active {
71        world_cache.active_cell_indices[compacted_index] = cell_id.x;
72    }
73
74    if thread_index == 1023u && workgroup_id.x == 1023u {
75        let active_cell_count = compacted_index + u32(cell_active);
76        world_cache.active_cells_count = active_cell_count;
77        world_cache_active_cells_dispatch = vec3((active_cell_count + 63u) / 64u, 1u, 1u);
78    }
79}