phosphor-crt 0.1.0

A real-time plotter of waveforms, imitating oscillscope CRTs
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

var<workgroup> workgroup_max: atomic<u32>;

@group(0) @binding(0) var intermediate_texture: texture_2d<f32>;
@group(0) @binding(1) var sampler_screen_linear: sampler;
@group(1) @binding(0) var<storage, read_write> global_max: atomic<u32>;

@compute @workgroup_size(16, 16, 1)
fn main(
    @builtin(global_invocation_id) global_invocation_id: vec3u,
    @builtin(local_invocation_id) local_invocation_id: vec3u,
) {
    if (local_invocation_id.x == 0 && local_invocation_id.y == 0) {
        atomicStore(&workgroup_max, u32(1));
    }
    workgroupBarrier();

    let texture_size = textureDimensions(intermediate_texture, 0);
    let position = global_invocation_id.xy;

    if (all(position < texture_size)) {
        let value: f32 = textureLoad(intermediate_texture, position, 0).x;
        // Scale the float to an integer range.
        // Since the texture is stored as an f16, the value is actually in
        // the range [0, 65504].
        let value_u32 = u32(value * f32(1 << 16));
        atomicMax(&workgroup_max, value_u32);
    }

    workgroupBarrier();

    if (local_invocation_id.x == 0 && local_invocation_id.y == 0) {
        atomicMax(&global_max, atomicLoad(&workgroup_max));
    }
}