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
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76

@group(0) @binding(0) var<uniform> params: StepParams;
@group(0) @binding(1) var<storage, read_write> body_states: array<BodyState>;
@group(0) @binding(2) var<storage, read> body_descs: array<BodyDescriptor>;
@group(0) @binding(3) var<storage, read_write> contacts: array<Contact>;
@group(0) @binding(4) var<storage, read> constraint_descs: array<ConstraintDescriptor>;
@group(0) @binding(5) var<storage, read_write> constraint_runtime: array<ConstraintRuntime>;
@group(0) @binding(6) var<storage, read_write> wake_flags: array<atomic<u32>>;
@group(0) @binding(7) var<storage, read> segments: array<u32>;
@group(0) @binding(8) var<storage, read> a_payload: array<u32>;
@group(0) @binding(9) var<storage, read_write> block_deltas: array<vec4f>;
@group(0) @binding(10) var<storage, read> block_counts: array<u32>;

struct BlockPair {
    first: Body,
    second: Body,
    split_first: Body,
    split_second: Body,
}

fn load_body(slot: u32) -> Body {
    return Body(body_states[slot], body_descs[slot]);
}

fn body_split(body: Body, scale: f32) -> Body {
    var split = body;
    split.desc.inverse_mass = body.desc.inverse_mass * scale;
    var inertia = body.desc.inverse_inertia;
    for (var index = 0u; index < 6u; index = index + 1u) {
        inertia[index] = inertia[index] * scale;
    }
    split.desc.inverse_inertia = inertia;
    return split;
}

fn block_scale(slot: u32) -> f32 {
    return f32(max(block_counts[slot], 1u));
}

fn block_pair(first_slot: u32, second_slot: u32) -> BlockPair {
    var pair: BlockPair;
    pair.first = load_body(first_slot);
    pair.second = load_body(second_slot);
    pair.split_first = body_split(pair.first, block_scale(first_slot));
    pair.split_second = body_split(pair.second, block_scale(second_slot));
    if (body_is_inert(pair.first)) {
        pair.first = body_frozen(pair.first);
        pair.split_first = body_frozen(pair.split_first);
    }
    if (body_is_inert(pair.second)) {
        pair.second = body_frozen(pair.second);
        pair.split_second = body_frozen(pair.split_second);
    }
    return pair;
}

fn store_block_delta(slot: u32, delta_a: vec3f, spin_a: vec3f, delta_b: vec3f, spin_b: vec3f) {
    block_deltas[slot * 4u] = vec4f(delta_a, 0.0);
    block_deltas[slot * 4u + 1u] = vec4f(spin_a, 0.0);
    block_deltas[slot * 4u + 2u] = vec4f(delta_b, 0.0);
    block_deltas[slot * 4u + 3u] = vec4f(spin_b, 0.0);
}

@compute @workgroup_size(WORKGROUP_SIZE)
fn main(@builtin(global_invocation_id) gid: vec3u) {
    let slot = gid.y * (WORKGROUPS_PER_ROW * WORKGROUP_SIZE) + gid.x;
    let contact_blocks = segments[SOLVER_BLOCK_CONTACT];
    if (slot >= contact_blocks + segments[SOLVER_BLOCK_CONSTRAINT]) {
        return;
    }
    let block = a_payload[slot];
    if (block < contact_blocks) {
        solve_contact_block(block, slot);
    } else {
        solve_constraint_block(block - contact_blocks, slot);
    }
}