wgpu-primitives 0.3.0

Composable GPU primitives for Rust applications using wgpu.
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

@group(0) @binding(0) var<storage, read_write> data: array<u32>;
@group(0) @binding(1) var<storage, read_write> aux: array<u32>;

const VT: u32 = {{VT}}u;
const BLOCK_SIZE: u32 = {{BLOCK_SIZE}}u;
const ITEMS_PER_BLOCK: u32 = VT * BLOCK_SIZE;

@compute @workgroup_size(BLOCK_SIZE)
fn main(
    @builtin(local_invocation_id) local_id: vec3<u32>,
    @builtin(workgroup_id) group_id: vec3<u32>
) {
    let tid = local_id.x;
    let flat_group_id = group_id.y * 65535u + group_id.x;
    
    if (flat_group_id == 0u) { return; }

    let group_base = flat_group_id * ITEMS_PER_BLOCK;
    let thread_base = group_base + (tid * VT);
    let add_val = aux[flat_group_id - 1u];

    for (var i = 0u; i < VT; i++) {
        let idx = thread_base + i;
        if (idx < arrayLength(&data)) {
            data[idx] += add_val;
        }
    }
}