1// Must have "common.wgsl" preprocessed before this file's contents.
2//
3// To compile this locally, you can run:
4// ```
5// cat common.wgsl timestamp_normalization.wgsl | cargo run -p naga-cli -- --stdin-file-path timestamp_normalization.wgsl
6// ```
7
8// For an explanation of the timestamp normalization process, see
9// the `mod.rs` file in this folder.
10
11// These is the timestamp period turned into a fraction
12// with an integer numerator and denominator. The denominator
13// is a power of two, so the division can be done with a shift.
14override TIMESTAMP_PERIOD_MULTIPLY: u32 = 1;
15override TIMESTAMP_PERIOD_SHIFT: u32 = 0;
16
17@group(0) @binding(0)
18var<storage, read_write> timestamps: array<Uint64>;
19
20struct ImmediateData {
21    timestamp_offset: u32,
22    timestamp_count: u32,
23}
24
25var<immediate> im: ImmediateData;
26
27@compute @workgroup_size(64)
28fn main(@builtin(global_invocation_id) id: vec3u) {
29    if id.x >= im.timestamp_count {
30        return;
31    }
32
33    let index = id.x + im.timestamp_offset;
34
35    let input_value = timestamps[index];
36
37    let tmp1 = u64_mul_u32(input_value, TIMESTAMP_PERIOD_MULTIPLY);
38    let tmp2 = shift_right_96(tmp1, TIMESTAMP_PERIOD_SHIFT);
39
40    timestamps[index] = truncate_u96_to_u64(tmp2);
41}