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) cdf_0 : Cdf_std430_0,
@align(8) init_volume_0 : f32,
@align(4) init_radius_0 : f32,
@align(8) mass_0 : f32,
};
@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
{
@align(8) 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
{
@align(4) 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
{
@align(8) 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)));
}
struct BodyVelocity_0
{
@align(8) linear_0 : vec2<f32>,
@align(8) angular_0 : f32,
};
fn body_velocity_at_point_0( center_of_mass_0 : vec2<f32>, vels_0 : BodyVelocity_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 : u32 = entryPointParams_hmap_entries_0[slot_0].state_0;
if((entryPointParams_hmap_entries_0[slot_0].state_0) == _S16)
{
var _S18 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S18;
}
else
{
if(_S17 == 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 : BlockVirtualId_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 : BlockVirtualId_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;
if((shared_nodes_0[_S43].particle_id_0) != u32(4294967295))
{
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;
shared_affine_0[_S43] = mat2x2<f32>(entryPointParams_particles_dyn_0[_S44].affine_0.data_0[i32(0)][i32(0)], entryPointParams_particles_dyn_0[_S44].affine_0.data_0[i32(0)][i32(1)], entryPointParams_particles_dyn_0[_S44].affine_0.data_0[i32(1)][i32(0)], entryPointParams_particles_dyn_0[_S44].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);
shared_nodes_0[_S43].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S44];
}
else
{
shared_affinities_0[_S43] = u32(0);
var _S45 : vec2<f32> = vec2<f32>(0.0f);
shared_normals_0[_S43] = _S45;
shared_pos_0[_S43].pt_0 = _S45;
shared_affine_0[_S43] = mat2x2<f32>(_S45, _S45);
shared_vel_mass_0[_S43] = vec3<f32>(0.0f);
}
k_2 = k_2 + u32(1);
}
j_2 = j_2 + u32(1);
}
i_6 = i_6 + u32(1);
}
return;
}
fn p2g_step_0( _S46 : u32, _S47 : f32, _S48 : u32, _S49 : u32) -> P2GStepResult_0
{
var impulse_2 : vec2<f32>;
var ang_impulse_2 : f32;
var new_momentum_velocity_mass_2 : vec3<f32>;
var _S50 : u32;
var _S51 : vec2<u32>;
var _S52 : vec2<f32>;
var _S53 : u32 = flatten_shared_shift_0(u32(2), u32(2));
var _S54 : vec3<f32> = vec3<f32>(0.0f);
var _S55 : vec2<f32> = vec2<f32>(0.0f);
for(;;)
{
for(;;)
{
for(;;)
{
var _S56 : u32 = _S46 - _S53;
_S50 = _S56;
var _S57 : u32 = _S56 + u32(22);
var _S58 : Position_0 = shared_pos_0[_S57];
var _S59 : mat2x2<f32> = shared_affine_0[_S57];
var _S60 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S57], _S47);
var _S61 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S60, _S47);
var _S62 : vec2<f32> = shared_vel_mass_0[_S57].xy;
var _S63 : f32 = shared_vel_mass_0[_S57].z;
const _S64 : vec2<u32> = vec2<u32>(u32(2), u32(2));
_S51 = _S64;
var _S65 : vec2<u32> = _S64 - vec2<u32>(u32(2), u32(2));
var _S66 : vec2<f32> = _S62 * vec2<f32>(_S63);
var _S67 : vec2<f32> = vec2<f32>(_S65);
var _S68 : vec2<f32> = vec2<f32>(_S47);
_S52 = _S68;
var _S69 : vec2<f32> = _S60 + _S67 * _S68;
var _S70 : f32 = _S61[i32(0)][_S65.x] * _S61[i32(1)][_S65.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S57]))
{
if(_S49 != u32(4294967295))
{
var _S71 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S72 : vec2<f32> = _S69 + _S58.pt_0;
var _S73 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S71, _S72);
var _S74 : vec2<f32> = (_S62 - (_S73 + project_velocity_0(_S62 - _S73, shared_normals_0[_S57]))) * vec2<f32>((_S70 * _S63));
var _S75 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S72;
var _S76 : f32 = dot(_S74, vec2<f32>(_S75.y, - _S75.x));
new_momentum_velocity_mass_2 = _S54;
ang_impulse_2 = _S76;
impulse_2 = _S74;
break;
}
new_momentum_velocity_mass_2 = _S54;
}
else
{
new_momentum_velocity_mass_2 = vec3<f32>((((_S59) * (_S69))) + _S66, _S63) * vec3<f32>(_S70);
}
ang_impulse_2 = 0.0f;
impulse_2 = _S55;
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S77 : u32 = _S50 + u32(2);
var _S78 : Position_0 = shared_pos_0[_S77];
var _S79 : mat2x2<f32> = shared_affine_0[_S77];
var _S80 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S77], _S47);
var _S81 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S80, _S47);
var _S82 : vec2<f32> = shared_vel_mass_0[_S77].xy;
var _S83 : f32 = shared_vel_mass_0[_S77].z;
var _S84 : vec2<u32> = _S51 - vec2<u32>(u32(2), u32(0));
var _S85 : vec2<f32> = _S82 * vec2<f32>(_S83);
var _S86 : vec2<f32> = _S80 + vec2<f32>(_S84) * _S52;
var _S87 : f32 = _S81[i32(0)][_S84.x] * _S81[i32(1)][_S84.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S77]))
{
if(_S49 != u32(4294967295))
{
var _S88 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S89 : vec2<f32> = _S86 + _S78.pt_0;
var _S90 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S88, _S89);
var _S91 : vec2<f32> = (_S82 - (_S90 + project_velocity_0(_S82 - _S90, shared_normals_0[_S77]))) * vec2<f32>((_S87 * _S83));
var _S92 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S89;
var impulse_3 : vec2<f32> = impulse_2 + _S91;
ang_impulse_2 = ang_impulse_2 + dot(_S91, vec2<f32>(_S92.y, - _S92.x));
impulse_2 = impulse_3;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S79) * (_S86))) + _S85, _S83) * vec3<f32>(_S87);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S93 : u32 = _S50 + u32(12);
var _S94 : Position_0 = shared_pos_0[_S93];
var _S95 : mat2x2<f32> = shared_affine_0[_S93];
var _S96 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S93], _S47);
var _S97 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S96, _S47);
var _S98 : vec2<f32> = shared_vel_mass_0[_S93].xy;
var _S99 : f32 = shared_vel_mass_0[_S93].z;
var _S100 : vec2<u32> = _S51 - vec2<u32>(u32(2), u32(1));
var _S101 : vec2<f32> = _S98 * vec2<f32>(_S99);
var _S102 : vec2<f32> = _S96 + vec2<f32>(_S100) * _S52;
var _S103 : f32 = _S97[i32(0)][_S100.x] * _S97[i32(1)][_S100.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S93]))
{
if(_S49 != u32(4294967295))
{
var _S104 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S105 : vec2<f32> = _S102 + _S94.pt_0;
var _S106 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S104, _S105);
var _S107 : vec2<f32> = (_S98 - (_S106 + project_velocity_0(_S98 - _S106, shared_normals_0[_S93]))) * vec2<f32>((_S103 * _S99));
var _S108 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S105;
var impulse_4 : vec2<f32> = impulse_2 + _S107;
ang_impulse_2 = ang_impulse_2 + dot(_S107, vec2<f32>(_S108.y, - _S108.x));
impulse_2 = impulse_4;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S95) * (_S102))) + _S101, _S99) * vec3<f32>(_S103);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S109 : u32 = _S50 + u32(20);
var _S110 : Position_0 = shared_pos_0[_S109];
var _S111 : mat2x2<f32> = shared_affine_0[_S109];
var _S112 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S109], _S47);
var _S113 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S112, _S47);
var _S114 : vec2<f32> = shared_vel_mass_0[_S109].xy;
var _S115 : f32 = shared_vel_mass_0[_S109].z;
var _S116 : vec2<u32> = _S51 - vec2<u32>(u32(0), u32(2));
var _S117 : vec2<f32> = _S114 * vec2<f32>(_S115);
var _S118 : vec2<f32> = _S112 + vec2<f32>(_S116) * _S52;
var _S119 : f32 = _S113[i32(0)][_S116.x] * _S113[i32(1)][_S116.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S109]))
{
if(_S49 != u32(4294967295))
{
var _S120 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S121 : vec2<f32> = _S118 + _S110.pt_0;
var _S122 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S120, _S121);
var _S123 : vec2<f32> = (_S114 - (_S122 + project_velocity_0(_S114 - _S122, shared_normals_0[_S109]))) * vec2<f32>((_S119 * _S115));
var _S124 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S121;
var impulse_5 : vec2<f32> = impulse_2 + _S123;
ang_impulse_2 = ang_impulse_2 + dot(_S123, vec2<f32>(_S124.y, - _S124.x));
impulse_2 = impulse_5;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S111) * (_S118))) + _S117, _S115) * vec3<f32>(_S119);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S125 : Position_0 = shared_pos_0[_S50];
var _S126 : mat2x2<f32> = shared_affine_0[_S50];
var _S127 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S50], _S47);
var _S128 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S127, _S47);
var _S129 : vec2<f32> = shared_vel_mass_0[_S50].xy;
var _S130 : f32 = shared_vel_mass_0[_S50].z;
var _S131 : vec2<u32> = _S51 - vec2<u32>(u32(0), u32(0));
var _S132 : vec2<f32> = _S129 * vec2<f32>(_S130);
var _S133 : vec2<f32> = _S127 + vec2<f32>(_S131) * _S52;
var _S134 : f32 = _S128[i32(0)][_S131.x] * _S128[i32(1)][_S131.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S50]))
{
if(_S49 != u32(4294967295))
{
var _S135 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S136 : vec2<f32> = _S133 + _S125.pt_0;
var _S137 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S135, _S136);
var _S138 : vec2<f32> = (_S129 - (_S137 + project_velocity_0(_S129 - _S137, shared_normals_0[_S50]))) * vec2<f32>((_S134 * _S130));
var _S139 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S136;
var impulse_6 : vec2<f32> = impulse_2 + _S138;
ang_impulse_2 = ang_impulse_2 + dot(_S138, vec2<f32>(_S139.y, - _S139.x));
impulse_2 = impulse_6;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S126) * (_S133))) + _S132, _S130) * vec3<f32>(_S134);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S140 : u32 = _S50 + u32(10);
var _S141 : Position_0 = shared_pos_0[_S140];
var _S142 : mat2x2<f32> = shared_affine_0[_S140];
var _S143 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S140], _S47);
var _S144 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S143, _S47);
var _S145 : vec2<f32> = shared_vel_mass_0[_S140].xy;
var _S146 : f32 = shared_vel_mass_0[_S140].z;
var _S147 : vec2<u32> = _S51 - vec2<u32>(u32(0), u32(1));
var _S148 : vec2<f32> = _S145 * vec2<f32>(_S146);
var _S149 : vec2<f32> = _S143 + vec2<f32>(_S147) * _S52;
var _S150 : f32 = _S144[i32(0)][_S147.x] * _S144[i32(1)][_S147.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S140]))
{
if(_S49 != u32(4294967295))
{
var _S151 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S152 : vec2<f32> = _S149 + _S141.pt_0;
var _S153 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S151, _S152);
var _S154 : vec2<f32> = (_S145 - (_S153 + project_velocity_0(_S145 - _S153, shared_normals_0[_S140]))) * vec2<f32>((_S150 * _S146));
var _S155 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S152;
var impulse_7 : vec2<f32> = impulse_2 + _S154;
ang_impulse_2 = ang_impulse_2 + dot(_S154, vec2<f32>(_S155.y, - _S155.x));
impulse_2 = impulse_7;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S142) * (_S149))) + _S148, _S146) * vec3<f32>(_S150);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S156 : u32 = _S50 + u32(21);
var _S157 : Position_0 = shared_pos_0[_S156];
var _S158 : mat2x2<f32> = shared_affine_0[_S156];
var _S159 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S156], _S47);
var _S160 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S159, _S47);
var _S161 : vec2<f32> = shared_vel_mass_0[_S156].xy;
var _S162 : f32 = shared_vel_mass_0[_S156].z;
var _S163 : vec2<u32> = _S51 - vec2<u32>(u32(1), u32(2));
var _S164 : vec2<f32> = _S161 * vec2<f32>(_S162);
var _S165 : vec2<f32> = _S159 + vec2<f32>(_S163) * _S52;
var _S166 : f32 = _S160[i32(0)][_S163.x] * _S160[i32(1)][_S163.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S156]))
{
if(_S49 != u32(4294967295))
{
var _S167 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S168 : vec2<f32> = _S165 + _S157.pt_0;
var _S169 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S167, _S168);
var _S170 : vec2<f32> = (_S161 - (_S169 + project_velocity_0(_S161 - _S169, shared_normals_0[_S156]))) * vec2<f32>((_S166 * _S162));
var _S171 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S168;
var impulse_8 : vec2<f32> = impulse_2 + _S170;
ang_impulse_2 = ang_impulse_2 + dot(_S170, vec2<f32>(_S171.y, - _S171.x));
impulse_2 = impulse_8;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S158) * (_S165))) + _S164, _S162) * vec3<f32>(_S166);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S172 : u32 = _S50 + u32(1);
var _S173 : Position_0 = shared_pos_0[_S172];
var _S174 : mat2x2<f32> = shared_affine_0[_S172];
var _S175 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S172], _S47);
var _S176 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S175, _S47);
var _S177 : vec2<f32> = shared_vel_mass_0[_S172].xy;
var _S178 : f32 = shared_vel_mass_0[_S172].z;
var _S179 : vec2<u32> = _S51 - vec2<u32>(u32(1), u32(0));
var _S180 : vec2<f32> = _S177 * vec2<f32>(_S178);
var _S181 : vec2<f32> = _S175 + vec2<f32>(_S179) * _S52;
var _S182 : f32 = _S176[i32(0)][_S179.x] * _S176[i32(1)][_S179.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S172]))
{
if(_S49 != u32(4294967295))
{
var _S183 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S184 : vec2<f32> = _S181 + _S173.pt_0;
var _S185 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S183, _S184);
var _S186 : vec2<f32> = (_S177 - (_S185 + project_velocity_0(_S177 - _S185, shared_normals_0[_S172]))) * vec2<f32>((_S182 * _S178));
var _S187 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S184;
var impulse_9 : vec2<f32> = impulse_2 + _S186;
ang_impulse_2 = ang_impulse_2 + dot(_S186, vec2<f32>(_S187.y, - _S187.x));
impulse_2 = impulse_9;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S174) * (_S181))) + _S180, _S178) * vec3<f32>(_S182);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S188 : u32 = _S50 + u32(11);
var _S189 : Position_0 = shared_pos_0[_S188];
var _S190 : mat2x2<f32> = shared_affine_0[_S188];
var _S191 : vec2<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S188], _S47);
var _S192 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S191, _S47);
var _S193 : vec2<f32> = shared_vel_mass_0[_S188].xy;
var _S194 : f32 = shared_vel_mass_0[_S188].z;
var _S195 : vec2<u32> = _S51 - vec2<u32>(u32(1), u32(1));
var _S196 : vec2<f32> = _S193 * vec2<f32>(_S194);
var _S197 : vec2<f32> = _S191 + vec2<f32>(_S195) * _S52;
var _S198 : f32 = _S192[i32(0)][_S195.x] * _S192[i32(1)][_S195.y];
if(!affinities_are_compatible_0(_S48, shared_affinities_0[_S188]))
{
if(_S49 != u32(4294967295))
{
var _S199 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S49].linear_0, entryPointParams_body_vels_0[_S49].angular_0 );
var _S200 : vec2<f32> = _S197 + _S189.pt_0;
var _S201 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S49].com_0, _S199, _S200);
var _S202 : vec2<f32> = (_S193 - (_S201 + project_velocity_0(_S193 - _S201, shared_normals_0[_S188]))) * vec2<f32>((_S198 * _S194));
var _S203 : vec2<f32> = entryPointParams_body_impulses_0[_S49].com_0 - _S200;
var impulse_10 : vec2<f32> = impulse_2 + _S202;
ang_impulse_2 = ang_impulse_2 + dot(_S202, vec2<f32>(_S203.y, - _S203.x));
impulse_2 = impulse_10;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec3<f32>((((_S190) * (_S197))) + _S196, _S194) * vec3<f32>(_S198);
}
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 _S204 : u32 = block_id_0.x;
var _S205 : BlockVirtualId_0 = BlockVirtualId_0( entryPointParams_active_blocks_0[_S204].virtual_id_0.id_0 );
if(tid_flat_0 == u32(0))
{
atomicStore(&(max_linked_list_length_0), u32(0));
}
workgroupBarrier();
fetch_max_linked_lists_length_0(tid_0, _S205, _S204);
workgroupBarrier();
var _S206 : u32 = atomicLoad(&(max_linked_list_length_0));
max_linked_list_length_uniform_0 = _S206;
fetch_nodes_0(tid_0, _S205, _S204);
var _S207 : u32 = flatten_shared_index_0(tid_0.x + u32(8), tid_0.y + u32(8));
var _S208 : u32 = shared_nodes_0[_S207].global_id_0;
var _S209 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S207].global_id_0].cdf_1.affinities_0;
var _S210 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S207].global_id_0].cdf_1.closest_id_0;
var total_result_0 : P2GStepResult_0 = P2GStepResult_x24init_0(vec3<f32>(0.0f, 0.0f, 0.0f), vec2<f32>(0.0f, 0.0f), 0.0f);
var _S211 : u32 = (workgroupUniformLoad(&((max_linked_list_length_uniform_0))));
var i_7 : u32 = u32(0);
for(;;)
{
if(i_7 < _S211)
{
}
else
{
break;
}
workgroupBarrier();
fetch_next_particle_0(tid_0);
workgroupBarrier();
var _S212 : P2GStepResult_0 = p2g_step_0(_S207, entryPointParams_grid_0[i32(0)].cell_width_0, _S209, _S210);
total_result_0.new_momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0 + _S212.new_momentum_velocity_mass_0;
total_result_0.impulse_0 = total_result_0.impulse_0 + _S212.impulse_0;
total_result_0.ang_impulse_0 = total_result_0.ang_impulse_0 + _S212.ang_impulse_0;
i_7 = i_7 + u32(1);
}
entryPointParams_nodes_0[_S208].momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0;
if(_S210 != u32(4294967295))
{
var _S213 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S210].linear_x_0), flt2int_0(total_result_0.impulse_0.x));
var _S214 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S210].linear_y_0), flt2int_0(total_result_0.impulse_0.y));
var _S215 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S210].angular_1), flt2int_0(total_result_0.ang_impulse_0));
}
return;
}