slosh2d 0.4.1

Cross-platform GPU 2D Material Point Method implementation.
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
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115

struct SimulationParams_std140_0
{
    @align(16) gravity_0 : vec2<f32>,
    @align(8) padding_0 : f32,
    @align(4) dt_0 : f32,
};

@binding(0) @group(0) var<uniform> entryPointParams_sim_params_0 : SimulationParams_std140_0;
struct GridGeneric_std430_0
{
    @align(4) num_active_blocks_0 : u32,
    @align(4) cell_width_0 : f32,
    @align(4) hmap_capacity_0 : u32,
    @align(4) capacity_0 : u32,
};

@binding(1) @group(0) var<storage, read> entryPointParams_grid_0 : array<GridGeneric_std430_0>;

struct BlockVirtualId_std430_0
{
    @align(8) id_0 : vec2<i32>,
};

struct ActiveBlockHeaderGeneric_std430_0
{
    @align(8) virtual_id_0 : BlockVirtualId_std430_0,
    @align(8) first_particle_0 : u32,
    @align(4) num_particles_0 : u32,
};

@binding(2) @group(0) var<storage, read> entryPointParams_active_blocks_0 : array<ActiveBlockHeaderGeneric_std430_0>;

struct NodeCdf_std430_0
{
    @align(4) distance_0 : f32,
    @align(4) affinities_0 : u32,
    @align(4) closest_id_0 : u32,
};

struct Node_std430_0
{
    @align(16) momentum_velocity_mass_0 : vec3<f32>,
    @align(4) cdf_0 : NodeCdf_std430_0,
};

@binding(3) @group(0) var<storage, read_write> entryPointParams_nodes_0 : array<Node_std430_0>;

struct BlockHeaderId_0
{
     id_1 : u32,
};

fn BlockHeaderId_x24init_0( i_0 : u32) -> BlockHeaderId_0
{
    var _S1 : BlockHeaderId_0;
    _S1.id_1 = i_0;
    return _S1;
}

struct BlockPhysicalId_0
{
     id_2 : u32,
};

fn BlockPhysicalId_x24init_0( i_1 : u32) -> BlockPhysicalId_0
{
    var _S2 : BlockPhysicalId_0;
    _S2.id_2 = i_1;
    return _S2;
}

fn block_header_id_to_physical_id_0( hid_0 : BlockHeaderId_0) -> BlockPhysicalId_0
{
    return BlockPhysicalId_x24init_0(hid_0.id_1 * u32(64));
}

struct NodePhysicalId_0
{
     id_3 : u32,
};

fn NodePhysicalId_x24init_0( i_2 : u32) -> NodePhysicalId_0
{
    var _S3 : NodePhysicalId_0;
    _S3.id_3 = i_2;
    return _S3;
}

fn node_id_0( pid_0 : BlockPhysicalId_0,  shift_in_block_0 : vec2<u32>) -> NodePhysicalId_0
{
    return NodePhysicalId_x24init_0(pid_0.id_2 + shift_in_block_0.x + shift_in_block_0.y * u32(8));
}

fn update_single_cell_0( sim_params_0 : ptr<function, SimulationParams_std140_0>,  cell_width_1 : f32,  cell_pos_0 : vec2<f32>,  momentum_velocity_mass_1 : vec3<f32>) -> vec3<f32>
{
    var _S4 : f32 = momentum_velocity_mass_1.z;
    var _S5 : vec2<f32> = vec2<f32>((cell_width_1 / (*sim_params_0).dt_0));
    return vec3<f32>(clamp((momentum_velocity_mass_1.xy + vec2<f32>(_S4) * (*sim_params_0).gravity_0 * vec2<f32>((*sim_params_0).dt_0)) * vec2<f32>(select(0.0f, 1.0f / _S4, _S4 > 0.0f)), (vec2<f32>(0) - _S5), _S5), _S4);
}

@compute
@workgroup_size(8, 8, 1)
fn grid_update(@builtin(workgroup_id) block_id_0 : vec3<u32>, @builtin(local_invocation_id) tid_0 : vec3<u32>, @builtin(local_invocation_index) tid_flat_0 : u32)
{
    var _S6 : u32 = block_id_0.x;
    var _S7 : GridGeneric_std430_0 = entryPointParams_grid_0[i32(0)];
    var _S8 : vec2<u32> = tid_0.xy;
    var _S9 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(BlockHeaderId_x24init_0(_S6)), _S8);
    var _S10 : vec2<f32> = vec2<f32>(entryPointParams_active_blocks_0[_S6].virtual_id_0.id_0 * vec2<i32>(i32(8)) + vec2<i32>(_S8)) * vec2<f32>(_S7.cell_width_0);
    var _S11 : SimulationParams_std140_0 = entryPointParams_sim_params_0;
    var _S12 : vec3<f32> = update_single_cell_0(&(_S11), _S7.cell_width_0, _S10, entryPointParams_nodes_0[_S9.id_3].momentum_velocity_mass_0);
    entryPointParams_nodes_0[_S9.id_3].momentum_velocity_mass_0 = _S12;
    return;
}