henad-models 0.3.0

Example models for Henad, a parallel agent-based modelling 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

// Leaf of the stat reduction. One workgroup folds its slice down to one value per lane, and
// `GpuLaneReduce` owns every level above this leaf.

#import henad::dispatch::WORKGROUP
#import henad::reduce_tree::block_sum
#import gpu_ants::state::HAS_FOOD_BIT

struct Params {
    n: u32,
    lanes: u32,
    groups_x: u32,
    num_agents: u32,
    n_cells: u32,
    _pad0: u32,
    _pad1: u32,
    _pad2: u32,
}

@group(0) @binding(0) var<storage, read> state: array<u32>;
@group(0) @binding(1) var<storage, read> field: array<f32>;
@group(0) @binding(2) var<storage, read_write> partials: array<f32>;
@group(0) @binding(3) var<uniform> params: Params;

@compute
@workgroup_size(256)
fn main(
    @builtin(local_invocation_id) lid: vec3<u32>,
    @builtin(workgroup_id) wid: vec3<u32>,
) {
    let block = wid.y * params.groups_x + wid.x;
    let i = block * WORKGROUP + lid.x;

    for (var lane: u32 = 0u; lane < params.lanes; lane = lane + 1u) {
        var value: f32 = 0.0;
        // Lane 0 is per ant and lane 1 per cell, so each lane bounds-checks its own domain.
        if (lane == 0u) {
            if (i < params.num_agents) {
                value = f32((state[i] & HAS_FOOD_BIT) != 0u);
            }
        } else {
            if (i < params.n_cells) {
                value = field[i] + field[params.n_cells + i];
            }
        }

        let total = block_sum(lid.x, value);
        if (lid.x == 0u) {
            partials[block * params.lanes + lane] = total;
        }
    }
}