vyre 0.4.0

GPU compute intermediate representation with a standard operation library
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

@group(0) @binding(0) var<storage, read> input_words: array<vec2<u32>>;
@group(0) @binding(1) var<storage, read_write> output_words: array<vec2<u32>>;

fn buffer_byte_swap_u64_part(value: u32) -> u32 {
    return ((value & 0x000000ffu) << 24u) |
        ((value & 0x0000ff00u) << 8u) |
        ((value >> 8u) & 0x0000ff00u) |
        ((value >> 24u) & 0x000000ffu);
}

fn buffer_byte_swap_u64_low(value: vec2<u32>) -> u32 {
    return buffer_byte_swap_u64_part(value.y);
}

fn buffer_byte_swap_u64_high(value: vec2<u32>) -> u32 {
    return buffer_byte_swap_u64_part(value.x);
}

fn buffer_byte_swap_u64_compose(value: vec2<u32>) -> vec2<u32> {
    let low = buffer_byte_swap_u64_low(value);
    let high = buffer_byte_swap_u64_high(value);
    return vec2<u32>(low, high);
}

fn buffer_byte_swap_u64_roundtrip_probe(value: vec2<u32>) -> vec2<u32> {
    let once = buffer_byte_swap_u64_compose(value);
    return buffer_byte_swap_u64_compose(once);
}

@compute @workgroup_size(64, 1, 1)
fn buffer_byte_swap_u64(@builtin(global_invocation_id) id: vec3<u32>) {
    let value = input_words[id.x];
    let swapped = buffer_byte_swap_u64_compose(value);
    let _probe = buffer_byte_swap_u64_roundtrip_probe(value);
    output_words[id.x] = swapped;
}