dynamis-world 0.6.0

GPU-driven physics engine
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

@group(0) @binding(0) var<storage, read> contacts: array<Contact>;
@group(0) @binding(1) var<storage, read_write> contact_count: array<atomic<u32>>;
@group(0) @binding(2) var<storage, read> body_states: array<BodyState>;
@group(0) @binding(3) var<storage, read_write> resting: array<Contact>;
@group(0) @binding(4) var<storage, read_write> resting_live: array<u32>;
@group(0) @binding(5) var<storage, read_write> resting_count: array<atomic<u32>>;
@group(0) @binding(6) var<storage, read_write> spillover: array<atomic<u32>>;
@group(0) @binding(7) var<storage, read> body_descs: array<BodyDescriptor>;
@group(0) @binding(8) var<storage, read_write> resting_next: array<u32>;
@group(0) @binding(9) var<storage, read_write> resting_free: array<atomic<u32>>;

fn acquire_slot() -> u32 {
    var slot = 0u;
    loop {
        let head = atomicLoad(&resting_free[0]);
        if (head == NO_SLOT) {
            slot = atomicAdd(&resting_count[0], 1u);
            break;
        }
        let next = resting_next[head];
        if (atomicCompareExchangeWeak(&resting_free[0], head, next).exchanged) {
            slot = head;
            break;
        }
    }
    return slot;
}

@compute @workgroup_size(WORKGROUP_SIZE)
fn main(@builtin(global_invocation_id) gid: vec3u) {
    let index = gid.y * (WORKGROUPS_PER_ROW * WORKGROUP_SIZE) + gid.x;
    if (index >= min(atomicLoad(&contact_count[0]), arrayLength(&contacts))) {
        return;
    }
    let contact = contacts[index];
    let first = contact.a / MAX_COLLIDERS_PER_BODY;
    let second = contact.b / MAX_COLLIDERS_PER_BODY;
    if (body_is_active(body_states[first], body_descs[first]) ||
        body_is_active(body_states[second], body_descs[second])) {
        return;
    }
    let slot = acquire_slot();
    if (slot >= arrayLength(&resting)) {
        atomicAdd(&spillover[0], 1u);
        return;
    }
    resting[slot] = contact;
    resting_live[slot] = 1u;
}