fidget-wgpu 0.5.0

WGPU backend for Fidget
Documentation 
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
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108

/// Merge all tile stages into a pixel array
@group(2) @binding(0) var<storage, read_write> tile64_zmin: array<atomic<u32>>;
@group(2) @binding(1) var<storage, read_write> tile16_zmin: array<atomic<u32>>;
@group(2) @binding(2) var<storage, read_write> tile4_zmin: array<atomic<u32>>;
@group(2) @binding(3) var<storage, read_write> voxels: array<u32>;

// Dispatched as an 2D workgroup across render_size pixels
@compute @workgroup_size(8, 8)
fn merge_main(
    @builtin(global_invocation_id) global_id: vec3u
) {
    // Clamp to render size (which is always >= image size)
    if global_id.x >= config.render_size.x ||
       global_id.y >= config.render_size.y
    {
        return;
    }

    // Precompute useful sizes
    let size64 = config.render_size / 64;
    let size16 = size64 * 4u;
    let size4 = size16 * 4u;

    // Compute indices within tile data (using render size bounds)
    let index64 = global_id.x / 64 + global_id.y / 64 * size64.x;
    let index16 = global_id.x / 16 + global_id.y / 16 * size16.x;
    let index4 = global_id.x / 4 + global_id.y / 4 * size4.x;
    let index1 = global_id.x + global_id.y * config.render_size.x;

    // Merge larger tiles into the voxels image
    var out = voxels[index1];
    out = max(out, atomicLoad(&tile64_zmin[index64]));
    out = max(out, atomicLoad(&tile16_zmin[index16]));
    out = max(out, atomicLoad(&tile4_zmin[index4]));

    // Clamp to image z size and write to the heightmap
    voxels[index1] = min(out, config.image_size.z);

    // Copying from high-res to low-res tiles is deliberately racey, because
    // it's to improve the odds of raymarching early-exit (but is not required
    // for correctness)

    // Copy from voxels to tile4_zmin
    if (global_id.x % 4) == 0 && (global_id.y % 4) == 0 {
        var all_present = true;
        var new_zmin = 0xFFFFFFFu;
        let corner = global_id.xy;
        for (var i=0u; i < 4u && all_present; i++) {
            for (var j=0u; j < 4u && all_present; j++) {
                let pos = corner + vec2u(i, j);
                let index = pos.x + pos.y * config.render_size.x;
                let v = voxels[index];
                if v != 0 {
                    new_zmin = min(new_zmin, v);
                } else {
                    all_present = false;
                }
            }
        }
        if all_present {
            atomicMax(&tile4_zmin[index4], new_zmin);
        }
    }

    // Copy from tile4_zmin to tile16_zmin
    if (global_id.x % 16) == 0 && (global_id.y % 16) == 0 {
        var all_present = true;
        var new_zmin = 0xFFFFFFFu;
        let corner = global_id.xy / 4;
        for (var i=0u; i < 4u && all_present; i++) {
            for (var j=0u; j < 4u && all_present; j++) {
                let pos = corner + vec2u(i, j);
                let index = pos.x + pos.y * size4.x;
                let v = atomicLoad(&tile4_zmin[index]);
                if v != 0 {
                    new_zmin = min(new_zmin, v);
                } else {
                    all_present = false;
                }
            }
        }
        if all_present {
            atomicMax(&tile16_zmin[index16], new_zmin);
        }
    }

    // Copy from tile16_zmin to tile64_zmin
    if (global_id.x % 64) == 0 && (global_id.y % 64) == 0 {
        var all_present = true;
        var new_zmin = 0xFFFFFFFu;
        let corner = global_id.xy / 16;
        for (var i=0u; i < 4u && all_present; i++) {
            for (var j=0u; j < 4u && all_present; j++) {
                let pos = corner + vec2u(i, j);
                let index = pos.x + pos.y * size16.x;
                let v = atomicLoad(&tile16_zmin[index]);
                if v != 0 {
                    new_zmin = min(new_zmin, v);
                } else {
                    all_present = false;
                }
            }
        }
        if all_present {
            atomicMax(&tile64_zmin[index64], new_zmin);
        }
    }
}