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(0) @group(0) var<storage, read> entryPointParams_grid_0 : array<GridGeneric_std430_0>;
struct BlockVirtualId_std430_0
{
@align(8) id_0 : vec2<i32>,
};
struct BlockHeaderId_std430_0
{
@align(4) id_1 : u32,
};
struct GridHashMapEntryGeneric_std430_0
{
@align(8) state_0 : u32,
@align(8) key_0 : BlockVirtualId_std430_0,
@align(8) value_0 : BlockHeaderId_std430_0,
};
@binding(1) @group(0) var<storage, read> entryPointParams_hmap_entries_0 : array<GridHashMapEntryGeneric_std430_0>;
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 NodeLinkedListGeneric_std430_0
{
@align(4) head_0 : u32,
@align(4) len_0 : u32,
};
@binding(3) @group(0) var<storage, read> entryPointParams_nodes_linked_lists_0 : array<NodeLinkedListGeneric_std430_0>;
@binding(4) @group(0) var<storage, read> entryPointParams_particle_node_linked_lists_0 : array<u32>;
struct Position_std430_0
{
@align(8) pt_0 : vec2<f32>,
};
@binding(5) @group(0) var<storage, read> entryPointParams_particles_pos_0 : array<Position_std430_0>;
struct _MatrixStorage_float2x2std430_0
{
@align(8) data_0 : array<vec2<f32>, i32(2)>,
};
struct Cdf_std430_0
{
@align(8) normal_0 : vec2<f32>,
@align(8) rigid_vel_0 : vec2<f32>,
@align(8) signed_distance_0 : f32,
@align(4) affinity_0 : u32,
};
struct Dynamics_std430_0
{
@align(8) velocity_0 : vec2<f32>,
@align(8) def_grad_0 : _MatrixStorage_float2x2std430_0,
@align(8) affine_0 : _MatrixStorage_float2x2std430_0,
@align(8) vel_grad_det_0 : f32,
@align(8) cdf_0 : Cdf_std430_0,
@align(8) init_volume_0 : f32,
@align(4) init_radius_0 : f32,
@align(8) mass_0 : f32,
@align(4) phase_0 : f32,
@align(8) enabled_0 : u32,
};
@binding(6) @group(0) var<storage, read> entryPointParams_particles_dyn_0 : array<Dynamics_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_1 : NodeCdf_std430_0,
};
@binding(7) @group(0) var<storage, read_write> entryPointParams_nodes_0 : array<Node_std430_0>;
struct BodyVelocity_std430_0
{
@align(8) linear_0 : vec2<f32>,
@align(8) angular_0 : f32,
};
@binding(8) @group(0) var<storage, read> entryPointParams_body_vels_0 : array<BodyVelocity_std430_0>;
struct IntegerImpulseAtomic_std430_0
{
@align(8) com_0 : vec2<f32>,
@align(8) linear_x_0 : atomic<i32>,
@align(4) linear_y_0 : atomic<i32>,
@align(8) angular_1 : atomic<i32>,
@align(4) padding_0 : i32,
};
@binding(9) @group(0) var<storage, read_write> entryPointParams_body_impulses_0 : array<IntegerImpulseAtomic_std430_0>;
var<workgroup> max_linked_list_length_0 : atomic<u32>;
struct BlockVirtualId_0
{
id_0 : vec2<i32>,
};
fn BlockVirtualId_x24init_0( i_0 : vec2<i32>) -> BlockVirtualId_0
{
var _S1 : BlockVirtualId_0;
_S1.id_0 = i_0;
return _S1;
}
fn pack_key_0( key_1 : BlockVirtualId_0) -> u32
{
return (((bitcast<u32>(key_1.id_0.x + i32(32767)) & (u32(65535)))) | (((((bitcast<u32>(key_1.id_0.y + i32(32767)) & (u32(65535)))) << (u32(16))))));
}
fn hash_0( packed_key_0 : u32) -> u32
{
var key_2 : u32 = packed_key_0 * u32(3432918353);
return ((((key_2 << (u32(15)))) | (((key_2 >> (u32(17))))))) * u32(461845907);
}
struct BlockHeaderId_0
{
id_1 : u32,
};
fn BlockHeaderId_x24init_0( i_1 : u32) -> BlockHeaderId_0
{
var _S2 : BlockHeaderId_0;
_S2.id_1 = i_1;
return _S2;
}
struct BlockPhysicalId_0
{
id_2 : u32,
};
fn BlockPhysicalId_x24init_0( i_2 : u32) -> BlockPhysicalId_0
{
var _S3 : BlockPhysicalId_0;
_S3.id_2 = i_2;
return _S3;
}
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_3 : u32) -> NodePhysicalId_0
{
var _S4 : NodePhysicalId_0;
_S4.id_3 = i_3;
return _S4;
}
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));
}
var<workgroup> max_linked_list_length_uniform_0 : u32;
fn flatten_shared_index_0( x_0 : u32, y_0 : u32) -> u32
{
return x_0 - u32(6) + (y_0 - u32(6)) * u32(10);
}
struct SharedNode_0
{
particle_id_0 : u32,
global_id_0 : u32,
};
var<workgroup> shared_nodes_0 : array<SharedNode_0, i32(100)>;
struct P2GStepResult_0
{
new_momentum_velocity_mass_0 : vec3<f32>,
impulse_0 : vec2<f32>,
ang_impulse_0 : f32,
};
fn P2GStepResult_x24init_0( new_momentum_velocity_mass_1 : vec3<f32>, impulse_1 : vec2<f32>, ang_impulse_1 : f32) -> P2GStepResult_0
{
var _S5 : P2GStepResult_0;
_S5.new_momentum_velocity_mass_0 = new_momentum_velocity_mass_1;
_S5.impulse_0 = impulse_1;
_S5.ang_impulse_0 = ang_impulse_1;
return _S5;
}
var<workgroup> shared_affinities_0 : array<u32, i32(100)>;
var<workgroup> shared_normals_0 : array<vec2<f32>, i32(100)>;
struct Position_0
{
pt_0 : vec2<f32>,
};
var<workgroup> shared_pos_0 : array<Position_0, i32(100)>;
var<workgroup> shared_affine_0 : array<mat2x2<f32>, i32(100)>;
var<workgroup> shared_vel_mass_0 : array<vec3<f32>, i32(100)>;
fn flatten_shared_shift_0( x_1 : u32, y_1 : u32) -> u32
{
return x_1 + y_1 * u32(10);
}
fn associated_grid_pos_0( part_pos_0 : Position_0, cell_width_1 : f32) -> vec2<f32>
{
var _S6 : vec2<f32> = vec2<f32>(cell_width_1);
return (round(part_pos_0.pt_0 / _S6) - vec2<f32>(1.0f)) * _S6;
}
fn dir_to_associated_grid_node_0( part_pos_1 : Position_0, cell_width_2 : f32) -> vec2<f32>
{
return associated_grid_pos_0(part_pos_1, cell_width_2) - part_pos_1.pt_0;
}
fn QuadraticKernel_eval_all_0( x_2 : f32) -> vec3<f32>
{
var _S7 : f32 = 1.5f - x_2;
var _S8 : f32 = x_2 - 1.0f;
var _S9 : f32 = x_2 - 0.5f;
return vec3<f32>(0.5f * _S7 * _S7, 0.75f - _S8 * _S8, 0.5f * _S9 * _S9);
}
fn QuadraticKernel_precompute_weights_0( ref_elt_pos_minus_particle_pos_0 : vec2<f32>, h_0 : f32) -> mat2x3<f32>
{
return mat2x3<f32>(QuadraticKernel_eval_all_0(- ref_elt_pos_minus_particle_pos_0.x / h_0), QuadraticKernel_eval_all_0(- ref_elt_pos_minus_particle_pos_0.y / h_0));
}
fn affinities_are_compatible_0( affinity1_0 : u32, affinity2_0 : u32) -> bool
{
var _S10 : u32 = (((affinity1_0 & (affinity2_0))) & (u32(65535)));
return ((((affinity1_0 >> (u32(16)))) & (_S10))) == ((((affinity2_0 >> (u32(16)))) & (_S10)));
}
fn body_velocity_at_point_0( center_of_mass_0 : vec2<f32>, vels_0 : ptr<function, BodyVelocity_std430_0>, point_0 : vec2<f32>) -> vec2<f32>
{
var _S11 : vec2<f32> = point_0 - center_of_mass_0;
return (*vels_0).linear_0 + vec2<f32>((*vels_0).angular_0) * vec2<f32>(- _S11.y, _S11.x);
}
fn project_velocity_0( vel_0 : vec2<f32>, n_0 : vec2<f32>) -> vec2<f32>
{
var _S12 : f32 = dot(vel_0, n_0);
if(_S12 < 0.0f)
{
var _S13 : vec2<f32> = vel_0 - n_0 * vec2<f32>(_S12);
var _S14 : f32 = length(_S13);
return select(vec2<f32>(0.0f), _S13 / vec2<f32>(_S14), _S14 > 9.99999993922529029e-09f) * vec2<f32>(max(0.0f, _S14 + 20.0f * _S12));
}
else
{
return vel_0;
}
}
fn flt2int_0( flt_0 : f32) -> i32
{
return i32(flt_0 * 1.0e+05f);
}
fn find_block_header_id_0( _S15 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S16 : u32 = pack_key_0(_S15);
var slot_0 : u32 = ((hash_0(_S16)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S17 : GridHashMapEntryGeneric_std430_0 = entryPointParams_hmap_entries_0[slot_0];
if((_S17.state_0) == _S16)
{
var _S18 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S18;
}
else
{
if((_S17.state_0) == u32(4294967295))
{
break;
}
}
slot_0 = ((slot_0 + u32(1)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
}
return BlockHeaderId_x24init_0(u32(4294967295));
}
fn fetch_max_linked_lists_length_0( _S19 : vec3<u32>, _S20 : ptr<function, BlockVirtualId_std430_0>, _S21 : u32)
{
var _S22 : vec2<i32> = (*_S20).id_0 - vec2<i32>(i32(1), i32(1));
var i_4 : u32 = u32(0);
for(;;)
{
if(i_4 <= u32(1))
{
}
else
{
break;
}
var j_0 : u32 = u32(0);
for(;;)
{
if(j_0 <= u32(1))
{
}
else
{
break;
}
var k_0 : u32 = u32(0);
for(;;)
{
if(k_0 <= u32(0))
{
}
else
{
break;
}
var _S23 : bool;
if(i_4 == u32(0))
{
_S23 = (_S19.x) < u32(6);
}
else
{
_S23 = false;
}
var _S24 : bool;
if(_S23)
{
_S24 = true;
}
else
{
if(j_0 == u32(0))
{
_S24 = (_S19.y) < u32(6);
}
else
{
_S24 = false;
}
}
if(_S24)
{
k_0 = k_0 + u32(1);
continue;
}
var _S25 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S22 + vec2<i32>(vec2<u32>(i_4, j_0))));
if((_S25.id_1) != u32(4294967295))
{
var _S26 : u32 = atomicMax(&(max_linked_list_length_0), entryPointParams_nodes_linked_lists_0[node_id_0(block_header_id_to_physical_id_0(_S25), _S19.xy).id_3].len_0);
}
k_0 = k_0 + u32(1);
}
j_0 = j_0 + u32(1);
}
i_4 = i_4 + u32(1);
}
return;
}
fn fetch_nodes_0( _S27 : vec3<u32>, _S28 : ptr<function, BlockVirtualId_std430_0>, _S29 : u32)
{
var _S30 : vec2<i32> = (*_S28).id_0 - vec2<i32>(i32(1), i32(1));
var i_5 : u32 = u32(0);
for(;;)
{
if(i_5 <= u32(1))
{
}
else
{
break;
}
var j_1 : u32 = u32(0);
for(;;)
{
if(j_1 <= u32(1))
{
}
else
{
break;
}
var k_1 : u32 = u32(0);
for(;;)
{
if(k_1 <= u32(0))
{
}
else
{
break;
}
var _S31 : bool;
if(i_5 == u32(0))
{
_S31 = (_S27.x) < u32(6);
}
else
{
_S31 = false;
}
var _S32 : bool;
if(_S31)
{
_S32 = true;
}
else
{
if(j_1 == u32(0))
{
_S32 = (_S27.y) < u32(6);
}
else
{
_S32 = false;
}
}
if(_S32)
{
k_1 = k_1 + u32(1);
continue;
}
var _S33 : vec2<u32> = vec2<u32>(i_5, j_1);
var _S34 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S30 + vec2<i32>(_S33)));
var _S35 : vec2<u32> = _S27.xy;
var _S36 : vec2<u32> = _S33 * vec2<u32>(u32(8)) + _S35;
var _S37 : u32 = flatten_shared_index_0(_S36.x, _S36.y);
if((_S34.id_1) != u32(4294967295))
{
var _S38 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S34), _S35);
shared_nodes_0[_S37].particle_id_0 = entryPointParams_nodes_linked_lists_0[_S38.id_3].head_0;
shared_nodes_0[_S37].global_id_0 = _S38.id_3;
}
else
{
shared_nodes_0[_S37].particle_id_0 = u32(4294967295);
}
k_1 = k_1 + u32(1);
}
j_1 = j_1 + u32(1);
}
i_5 = i_5 + u32(1);
}
return;
}
fn fetch_next_particle_0( _S39 : vec3<u32>)
{
var i_6 : u32 = u32(0);
for(;;)
{
if(i_6 <= u32(1))
{
}
else
{
break;
}
var j_2 : u32 = u32(0);
for(;;)
{
if(j_2 <= u32(1))
{
}
else
{
break;
}
var k_2 : u32 = u32(0);
for(;;)
{
if(k_2 <= u32(0))
{
}
else
{
break;
}
var _S40 : bool;
if(i_6 == u32(0))
{
_S40 = (_S39.x) < u32(6);
}
else
{
_S40 = false;
}
var _S41 : bool;
if(_S40)
{
_S41 = true;
}
else
{
if(j_2 == u32(0))
{
_S41 = (_S39.y) < u32(6);
}
else
{
_S41 = false;
}
}
if(_S41)
{
k_2 = k_2 + u32(1);
continue;
}
var _S42 : vec2<u32> = vec2<u32>(i_6, j_2) * vec2<u32>(u32(8)) + _S39.xy;
var _S43 : u32 = flatten_shared_index_0(_S42.x, _S42.y);
var _S44 : u32 = shared_nodes_0[_S43].particle_id_0;
var _S45 : bool = (shared_nodes_0[_S43].particle_id_0) != u32(4294967295);
var _S46 : bool;
if(_S45)
{
_S46 = (entryPointParams_particles_dyn_0[_S44].enabled_0) != u32(0);
}
else
{
_S46 = false;
}
if(_S46)
{
shared_affinities_0[_S43] = entryPointParams_particles_dyn_0[_S44].cdf_0.affinity_0;
shared_normals_0[_S43] = entryPointParams_particles_dyn_0[_S44].cdf_0.normal_0;
shared_pos_0[_S43].pt_0 = entryPointParams_particles_pos_0[_S44].pt_0;
var _S47 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S44];
shared_affine_0[_S43] = mat2x2<f32>(_S47.affine_0.data_0[i32(0)][i32(0)], _S47.affine_0.data_0[i32(0)][i32(1)], _S47.affine_0.data_0[i32(1)][i32(0)], _S47.affine_0.data_0[i32(1)][i32(1)]);
shared_vel_mass_0[_S43] = vec3<f32>(entryPointParams_particles_dyn_0[_S44].velocity_0, entryPointParams_particles_dyn_0[_S44].mass_0);
}
else
{
shared_affinities_0[_S43] = u32(0);
var _S48 : vec2<f32> = vec2<f32>(0.0f);
shared_normals_0[_S43] = _S48;
shared_pos_0[_S43].pt_0 = _S48;
shared_affine_0[_S43] = mat2x2<f32>(_S48, _S48);
shared_vel_mass_0[_S43] = vec3<f32>(0.0f);
}
if(_S45)
{
shared_nodes_0[_S43].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S44];
}
k_2 = k_2 + u32(1);
}
j_2 = j_2 + u32(1);
}
i_6 = i_6 + u32(1);
}
return;
}
fn p2g_step_0( _S49 : u32, _S50 : f32, _S51 : u32, _S52 : u32) -> P2GStepResult_0
{
var impulse_2 : vec2<f32>;
var ang_impulse_2 : f32;
var new_momentum_velocity_mass_2 : vec3<f32>;
var _S53 : u32;
var _S54 : vec2<u32>;
var _S55 : vec2<f32>;
var _S56 : u32 = flatten_shared_shift_0(u32(2), u32(2));
var _S57 : vec3<f32> = vec3<f32>(0.0f);
var _S58 : vec2<f32> = vec2<f32>(0.0f);
for(;;)
{
for(;;)
{
for(;;)
{
var _S59 : u32 = _S49 - _S56;
_S53 = _S59;
var _S60 : u32 = _S59 + u32(22);
var _S61 : Position_0 = shared_pos_0[_S60];
var _S62 : mat2x2<f32> = shared_affine_0[_S60];
var _S63 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S60], _S50);
var _S64 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S63, _S50);
var _S65 : vec2<f32> = shared_vel_mass_0[_S60].xy;
var _S66 : f32 = shared_vel_mass_0[_S60].z;
const _S67 : vec2<u32> = vec2<u32>(u32(2), u32(2));
_S54 = _S67;
var _S68 : vec2<u32> = _S67 - vec2<u32>(u32(2), u32(2));
var _S69 : vec2<f32> = _S65 * vec2<f32>(_S66);
var _S70 : vec2<f32> = vec2<f32>(_S68);
var _S71 : vec2<f32> = vec2<f32>(_S50);
_S55 = _S71;
var _S72 : vec2<f32> = _S63 + _S70 * _S71;
var _S73 : f32 = _S64[i32(0)][_S68.x] * _S64[i32(1)][_S68.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S60]))
{
if(_S52 != u32(4294967295))
{
var _S74 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S75 : vec2<f32> = _S72 + _S61.pt_0;
var _S76 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S74), _S75);
var _S77 : vec2<f32> = (_S65 - (_S76 + project_velocity_0(_S65 - _S76, shared_normals_0[_S60]))) * vec2<f32>((_S73 * _S66));
var _S78 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S75;
var _S79 : f32 = dot(_S77, vec2<f32>(_S78.y, - _S78.x));
new_momentum_velocity_mass_2 = _S57;
ang_impulse_2 = _S79;
impulse_2 = _S77;
break;
}
new_momentum_velocity_mass_2 = _S57;
}
else
{
new_momentum_velocity_mass_2 = vec3<f32>((((_S62) * (_S72))) + _S69, _S66) * vec3<f32>(_S73);
}
ang_impulse_2 = 0.0f;
impulse_2 = _S58;
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S80 : u32 = _S53 + u32(2);
var _S81 : Position_0 = shared_pos_0[_S80];
var _S82 : mat2x2<f32> = shared_affine_0[_S80];
var _S83 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S80], _S50);
var _S84 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S83, _S50);
var _S85 : vec2<f32> = shared_vel_mass_0[_S80].xy;
var _S86 : f32 = shared_vel_mass_0[_S80].z;
var _S87 : vec2<u32> = _S54 - vec2<u32>(u32(2), u32(0));
var _S88 : vec2<f32> = _S85 * vec2<f32>(_S86);
var _S89 : vec2<f32> = _S83 + vec2<f32>(_S87) * _S55;
var _S90 : f32 = _S84[i32(0)][_S87.x] * _S84[i32(1)][_S87.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S80]))
{
if(_S52 != u32(4294967295))
{
var _S91 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S92 : vec2<f32> = _S89 + _S81.pt_0;
var _S93 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S91), _S92);
var _S94 : vec2<f32> = (_S85 - (_S93 + project_velocity_0(_S85 - _S93, shared_normals_0[_S80]))) * vec2<f32>((_S90 * _S86));
var _S95 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S92;
var impulse_3 : vec2<f32> = impulse_2 + _S94;
ang_impulse_2 = ang_impulse_2 + dot(_S94, vec2<f32>(_S95.y, - _S95.x));
impulse_2 = impulse_3;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S82) * (_S89))) + _S88, _S86) * vec3<f32>(_S90);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S96 : u32 = _S53 + u32(12);
var _S97 : Position_0 = shared_pos_0[_S96];
var _S98 : mat2x2<f32> = shared_affine_0[_S96];
var _S99 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S96], _S50);
var _S100 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S99, _S50);
var _S101 : vec2<f32> = shared_vel_mass_0[_S96].xy;
var _S102 : f32 = shared_vel_mass_0[_S96].z;
var _S103 : vec2<u32> = _S54 - vec2<u32>(u32(2), u32(1));
var _S104 : vec2<f32> = _S101 * vec2<f32>(_S102);
var _S105 : vec2<f32> = _S99 + vec2<f32>(_S103) * _S55;
var _S106 : f32 = _S100[i32(0)][_S103.x] * _S100[i32(1)][_S103.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S96]))
{
if(_S52 != u32(4294967295))
{
var _S107 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S108 : vec2<f32> = _S105 + _S97.pt_0;
var _S109 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S107), _S108);
var _S110 : vec2<f32> = (_S101 - (_S109 + project_velocity_0(_S101 - _S109, shared_normals_0[_S96]))) * vec2<f32>((_S106 * _S102));
var _S111 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S108;
var impulse_4 : vec2<f32> = impulse_2 + _S110;
ang_impulse_2 = ang_impulse_2 + dot(_S110, vec2<f32>(_S111.y, - _S111.x));
impulse_2 = impulse_4;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S98) * (_S105))) + _S104, _S102) * vec3<f32>(_S106);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S112 : u32 = _S53 + u32(20);
var _S113 : Position_0 = shared_pos_0[_S112];
var _S114 : mat2x2<f32> = shared_affine_0[_S112];
var _S115 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S112], _S50);
var _S116 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S115, _S50);
var _S117 : vec2<f32> = shared_vel_mass_0[_S112].xy;
var _S118 : f32 = shared_vel_mass_0[_S112].z;
var _S119 : vec2<u32> = _S54 - vec2<u32>(u32(0), u32(2));
var _S120 : vec2<f32> = _S117 * vec2<f32>(_S118);
var _S121 : vec2<f32> = _S115 + vec2<f32>(_S119) * _S55;
var _S122 : f32 = _S116[i32(0)][_S119.x] * _S116[i32(1)][_S119.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S112]))
{
if(_S52 != u32(4294967295))
{
var _S123 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S124 : vec2<f32> = _S121 + _S113.pt_0;
var _S125 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S123), _S124);
var _S126 : vec2<f32> = (_S117 - (_S125 + project_velocity_0(_S117 - _S125, shared_normals_0[_S112]))) * vec2<f32>((_S122 * _S118));
var _S127 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S124;
var impulse_5 : vec2<f32> = impulse_2 + _S126;
ang_impulse_2 = ang_impulse_2 + dot(_S126, vec2<f32>(_S127.y, - _S127.x));
impulse_2 = impulse_5;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S114) * (_S121))) + _S120, _S118) * vec3<f32>(_S122);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S128 : Position_0 = shared_pos_0[_S53];
var _S129 : mat2x2<f32> = shared_affine_0[_S53];
var _S130 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S53], _S50);
var _S131 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S130, _S50);
var _S132 : vec2<f32> = shared_vel_mass_0[_S53].xy;
var _S133 : f32 = shared_vel_mass_0[_S53].z;
var _S134 : vec2<u32> = _S54 - vec2<u32>(u32(0), u32(0));
var _S135 : vec2<f32> = _S132 * vec2<f32>(_S133);
var _S136 : vec2<f32> = _S130 + vec2<f32>(_S134) * _S55;
var _S137 : f32 = _S131[i32(0)][_S134.x] * _S131[i32(1)][_S134.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S53]))
{
if(_S52 != u32(4294967295))
{
var _S138 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S139 : vec2<f32> = _S136 + _S128.pt_0;
var _S140 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S138), _S139);
var _S141 : vec2<f32> = (_S132 - (_S140 + project_velocity_0(_S132 - _S140, shared_normals_0[_S53]))) * vec2<f32>((_S137 * _S133));
var _S142 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S139;
var impulse_6 : vec2<f32> = impulse_2 + _S141;
ang_impulse_2 = ang_impulse_2 + dot(_S141, vec2<f32>(_S142.y, - _S142.x));
impulse_2 = impulse_6;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S129) * (_S136))) + _S135, _S133) * vec3<f32>(_S137);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S143 : u32 = _S53 + u32(10);
var _S144 : Position_0 = shared_pos_0[_S143];
var _S145 : mat2x2<f32> = shared_affine_0[_S143];
var _S146 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S143], _S50);
var _S147 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S146, _S50);
var _S148 : vec2<f32> = shared_vel_mass_0[_S143].xy;
var _S149 : f32 = shared_vel_mass_0[_S143].z;
var _S150 : vec2<u32> = _S54 - vec2<u32>(u32(0), u32(1));
var _S151 : vec2<f32> = _S148 * vec2<f32>(_S149);
var _S152 : vec2<f32> = _S146 + vec2<f32>(_S150) * _S55;
var _S153 : f32 = _S147[i32(0)][_S150.x] * _S147[i32(1)][_S150.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S143]))
{
if(_S52 != u32(4294967295))
{
var _S154 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S155 : vec2<f32> = _S152 + _S144.pt_0;
var _S156 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S154), _S155);
var _S157 : vec2<f32> = (_S148 - (_S156 + project_velocity_0(_S148 - _S156, shared_normals_0[_S143]))) * vec2<f32>((_S153 * _S149));
var _S158 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S155;
var impulse_7 : vec2<f32> = impulse_2 + _S157;
ang_impulse_2 = ang_impulse_2 + dot(_S157, vec2<f32>(_S158.y, - _S158.x));
impulse_2 = impulse_7;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S145) * (_S152))) + _S151, _S149) * vec3<f32>(_S153);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S159 : u32 = _S53 + u32(21);
var _S160 : Position_0 = shared_pos_0[_S159];
var _S161 : mat2x2<f32> = shared_affine_0[_S159];
var _S162 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S159], _S50);
var _S163 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S162, _S50);
var _S164 : vec2<f32> = shared_vel_mass_0[_S159].xy;
var _S165 : f32 = shared_vel_mass_0[_S159].z;
var _S166 : vec2<u32> = _S54 - vec2<u32>(u32(1), u32(2));
var _S167 : vec2<f32> = _S164 * vec2<f32>(_S165);
var _S168 : vec2<f32> = _S162 + vec2<f32>(_S166) * _S55;
var _S169 : f32 = _S163[i32(0)][_S166.x] * _S163[i32(1)][_S166.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S159]))
{
if(_S52 != u32(4294967295))
{
var _S170 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S171 : vec2<f32> = _S168 + _S160.pt_0;
var _S172 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S170), _S171);
var _S173 : vec2<f32> = (_S164 - (_S172 + project_velocity_0(_S164 - _S172, shared_normals_0[_S159]))) * vec2<f32>((_S169 * _S165));
var _S174 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S171;
var impulse_8 : vec2<f32> = impulse_2 + _S173;
ang_impulse_2 = ang_impulse_2 + dot(_S173, vec2<f32>(_S174.y, - _S174.x));
impulse_2 = impulse_8;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S161) * (_S168))) + _S167, _S165) * vec3<f32>(_S169);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S175 : u32 = _S53 + u32(1);
var _S176 : Position_0 = shared_pos_0[_S175];
var _S177 : mat2x2<f32> = shared_affine_0[_S175];
var _S178 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S175], _S50);
var _S179 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S178, _S50);
var _S180 : vec2<f32> = shared_vel_mass_0[_S175].xy;
var _S181 : f32 = shared_vel_mass_0[_S175].z;
var _S182 : vec2<u32> = _S54 - vec2<u32>(u32(1), u32(0));
var _S183 : vec2<f32> = _S180 * vec2<f32>(_S181);
var _S184 : vec2<f32> = _S178 + vec2<f32>(_S182) * _S55;
var _S185 : f32 = _S179[i32(0)][_S182.x] * _S179[i32(1)][_S182.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S175]))
{
if(_S52 != u32(4294967295))
{
var _S186 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S187 : vec2<f32> = _S184 + _S176.pt_0;
var _S188 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S186), _S187);
var _S189 : vec2<f32> = (_S180 - (_S188 + project_velocity_0(_S180 - _S188, shared_normals_0[_S175]))) * vec2<f32>((_S185 * _S181));
var _S190 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S187;
var impulse_9 : vec2<f32> = impulse_2 + _S189;
ang_impulse_2 = ang_impulse_2 + dot(_S189, vec2<f32>(_S190.y, - _S190.x));
impulse_2 = impulse_9;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S177) * (_S184))) + _S183, _S181) * vec3<f32>(_S185);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S191 : u32 = _S53 + u32(11);
var _S192 : Position_0 = shared_pos_0[_S191];
var _S193 : mat2x2<f32> = shared_affine_0[_S191];
var _S194 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S191], _S50);
var _S195 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S194, _S50);
var _S196 : vec2<f32> = shared_vel_mass_0[_S191].xy;
var _S197 : f32 = shared_vel_mass_0[_S191].z;
var _S198 : vec2<u32> = _S54 - vec2<u32>(u32(1), u32(1));
var _S199 : vec2<f32> = _S196 * vec2<f32>(_S197);
var _S200 : vec2<f32> = _S194 + vec2<f32>(_S198) * _S55;
var _S201 : f32 = _S195[i32(0)][_S198.x] * _S195[i32(1)][_S198.y];
if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S191]))
{
if(_S52 != u32(4294967295))
{
var _S202 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
var _S203 : vec2<f32> = _S200 + _S192.pt_0;
var _S204 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S202), _S203);
var _S205 : vec2<f32> = (_S196 - (_S204 + project_velocity_0(_S196 - _S204, shared_normals_0[_S191]))) * vec2<f32>((_S201 * _S197));
var _S206 : vec2<f32> = entryPointParams_body_impulses_0[_S52].com_0 - _S203;
var impulse_10 : vec2<f32> = impulse_2 + _S205;
ang_impulse_2 = ang_impulse_2 + dot(_S205, vec2<f32>(_S206.y, - _S206.x));
impulse_2 = impulse_10;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S193) * (_S200))) + _S199, _S197) * vec3<f32>(_S201);
}
break;
}
break;
}
break;
}
return P2GStepResult_x24init_0(new_momentum_velocity_mass_2, impulse_2, ang_impulse_2);
}
@compute
@workgroup_size(8, 8, 1)
fn p2g(@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 _S207 : u32 = block_id_0.x;
var _S208 : ActiveBlockHeaderGeneric_std430_0 = entryPointParams_active_blocks_0[_S207];
if(tid_flat_0 == u32(0))
{
atomicStore(&(max_linked_list_length_0), u32(0));
}
workgroupBarrier();
fetch_max_linked_lists_length_0(tid_0, &(_S208.virtual_id_0), _S207);
workgroupBarrier();
var _S209 : u32 = atomicLoad(&(max_linked_list_length_0));
max_linked_list_length_uniform_0 = _S209;
fetch_nodes_0(tid_0, &(_S208.virtual_id_0), _S207);
var _S210 : u32 = flatten_shared_index_0(tid_0.x + u32(8), tid_0.y + u32(8));
var _S211 : u32 = shared_nodes_0[_S210].global_id_0;
var _S212 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S210].global_id_0].cdf_1.affinities_0;
var _S213 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S210].global_id_0].cdf_1.closest_id_0;
var total_result_0 : P2GStepResult_0 = P2GStepResult_x24init_0(vec3<f32>(0.0f), vec2<f32>(0.0f), 0.0f);
var _S214 : u32 = (workgroupUniformLoad(&((max_linked_list_length_uniform_0))));
var i_7 : u32 = u32(0);
for(;;)
{
if(i_7 < _S214)
{
}
else
{
break;
}
workgroupBarrier();
fetch_next_particle_0(tid_0);
workgroupBarrier();
var _S215 : P2GStepResult_0 = p2g_step_0(_S210, entryPointParams_grid_0[i32(0)].cell_width_0, _S212, _S213);
total_result_0.new_momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0 + _S215.new_momentum_velocity_mass_0;
total_result_0.impulse_0 = total_result_0.impulse_0 + _S215.impulse_0;
total_result_0.ang_impulse_0 = total_result_0.ang_impulse_0 + _S215.ang_impulse_0;
i_7 = i_7 + u32(1);
}
entryPointParams_nodes_0[_S211].momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0;
if(_S213 != u32(4294967295))
{
var _S216 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S213].linear_x_0), flt2int_0(total_result_0.impulse_0.x));
var _S217 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S213].linear_y_0), flt2int_0(total_result_0.impulse_0.y));
var _S218 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S213].angular_1), flt2int_0(total_result_0.ang_impulse_0));
}
return;
}