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(16) id_0 : vec3<i32>,
};
struct BlockHeaderId_std430_0
{
@align(4) id_1 : u32,
};
struct GridHashMapEntryGeneric_std430_0
{
@align(16) state_0 : u32,
@align(16) key_0 : BlockVirtualId_std430_0,
@align(16) 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(16) virtual_id_0 : BlockVirtualId_std430_0,
@align(16) 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(16) pt_0 : vec3<f32>,
};
@binding(5) @group(0) var<storage, read> entryPointParams_particles_pos_0 : array<Position_std430_0>;
struct _MatrixStorage_float3x3std430_0
{
@align(16) data_0 : array<vec3<f32>, i32(3)>,
};
struct Cdf_std430_0
{
@align(16) normal_0 : vec3<f32>,
@align(16) rigid_vel_0 : vec3<f32>,
@align(4) signed_distance_0 : f32,
@align(16) affinity_0 : u32,
};
struct Dynamics_std430_0
{
@align(16) velocity_0 : vec3<f32>,
@align(16) def_grad_0 : _MatrixStorage_float3x3std430_0,
@align(16) affine_0 : _MatrixStorage_float3x3std430_0,
@align(16) cdf_0 : Cdf_std430_0,
@align(16) 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 : vec4<f32>,
@align(16) 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(16) linear_0 : vec3<f32>,
@align(16) angular_0 : vec3<f32>,
};
@binding(8) @group(0) var<storage, read> entryPointParams_body_vels_0 : array<BodyVelocity_std430_0>;
struct IntegerImpulseAtomic_std430_0
{
@align(16) com_0 : vec3<f32>,
@align(4) linear_x_0 : atomic<i32>,
@align(16) linear_y_0 : atomic<i32>,
@align(4) linear_z_0 : atomic<i32>,
@align(8) padding_a_0 : i32,
@align(4) angular_x_0 : atomic<i32>,
@align(16) angular_y_0 : atomic<i32>,
@align(4) angular_z_0 : atomic<i32>,
@align(8) padding_b_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(16) id_0 : vec3<i32>,
};
fn BlockVirtualId_x24init_0( i_0 : vec3<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(1023)) & (u32(2047)))) | (((((bitcast<u32>(key_1.id_0.y + i32(511)) & (u32(1023)))) << (u32(11))))))) | (((((bitcast<u32>(key_1.id_0.z + i32(1023)) & (u32(2047)))) << (u32(21))))));
}
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 : vec3<u32>) -> NodePhysicalId_0
{
return NodePhysicalId_x24init_0(pid_0.id_2 + shift_in_block_0.x + shift_in_block_0.y * u32(4) + shift_in_block_0.z * u32(4) * u32(4));
}
var<workgroup> max_linked_list_length_uniform_0 : u32;
fn flatten_shared_index_0( x_0 : u32, y_0 : u32, z_0 : u32) -> u32
{
return x_0 - u32(2) + (y_0 - u32(2)) * u32(6) + (z_0 - u32(2)) * u32(6) * u32(6);
}
struct SharedNode_0
{
particle_id_0 : u32,
global_id_0 : u32,
};
var<workgroup> shared_nodes_0 : array<SharedNode_0, i32(216)>;
struct P2GStepResult_0
{
new_momentum_velocity_mass_0 : vec4<f32>,
impulse_0 : vec3<f32>,
ang_impulse_0 : vec3<f32>,
};
fn P2GStepResult_x24init_0( new_momentum_velocity_mass_1 : vec4<f32>, impulse_1 : vec3<f32>, ang_impulse_1 : vec3<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(216)>;
var<workgroup> shared_normals_0 : array<vec3<f32>, i32(216)>;
struct Position_0
{
@align(16) pt_0 : vec3<f32>,
};
var<workgroup> shared_pos_0 : array<Position_0, i32(216)>;
var<workgroup> shared_affine_0 : array<mat3x3<f32>, i32(216)>;
var<workgroup> shared_vel_mass_0 : array<vec4<f32>, i32(216)>;
fn flatten_shared_shift_0( x_1 : u32, y_1 : u32, z_1 : u32) -> u32
{
return x_1 + y_1 * u32(6) + z_1 * u32(6) * u32(6);
}
fn associated_grid_pos_0( part_pos_0 : Position_0, cell_width_1 : f32) -> vec3<f32>
{
var _S6 : vec3<f32> = vec3<f32>(cell_width_1);
return (round(part_pos_0.pt_0 / _S6) - vec3<f32>(1.0f)) * _S6;
}
fn dir_to_associated_grid_node_0( part_pos_1 : Position_0, cell_width_2 : f32) -> vec3<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 : vec3<f32>, h_0 : f32) -> mat3x3<f32>
{
return mat3x3<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), QuadraticKernel_eval_all_0(- ref_elt_pos_minus_particle_pos_0.z / 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(16) linear_0 : vec3<f32>,
@align(16) angular_0 : vec3<f32>,
};
fn body_velocity_at_point_0( com_1 : vec3<f32>, vels_0 : BodyVelocity_0, point_0 : vec3<f32>) -> vec3<f32>
{
return vels_0.linear_0 + cross(vels_0.angular_0, point_0 - com_1);
}
fn project_velocity_0( vel_0 : vec3<f32>, n_0 : vec3<f32>) -> vec3<f32>
{
var _S11 : f32 = dot(vel_0, n_0);
if(_S11 < 0.0f)
{
var _S12 : vec3<f32> = vel_0 - n_0 * vec3<f32>(_S11);
var _S13 : f32 = length(_S12);
return select(vec3<f32>(0.0f), _S12 / vec3<f32>(_S13), _S13 > 9.99999993922529029e-09f) * vec3<f32>(max(0.0f, _S13 + 20.0f * _S11));
}
else
{
return vel_0;
}
}
fn flt2int_0( flt_0 : f32) -> i32
{
return i32(flt_0 * 1.0e+05f);
}
fn find_block_header_id_0( _S14 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S15 : u32 = pack_key_0(_S14);
var slot_0 : u32 = ((hash_0(_S15)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S16 : u32 = entryPointParams_hmap_entries_0[slot_0].state_0;
if((entryPointParams_hmap_entries_0[slot_0].state_0) == _S15)
{
var _S17 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S17;
}
else
{
if(_S16 == 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( _S18 : vec3<u32>, _S19 : BlockVirtualId_0, _S20 : u32)
{
var _S21 : vec3<i32> = _S19.id_0 - vec3<i32>(i32(1), 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(1))
{
}
else
{
break;
}
var _S22 : bool;
if(i_4 == u32(0))
{
_S22 = (_S18.x) < u32(2);
}
else
{
_S22 = false;
}
var _S23 : bool;
if(_S22)
{
_S23 = true;
}
else
{
if(j_0 == u32(0))
{
_S23 = (_S18.y) < u32(2);
}
else
{
_S23 = false;
}
}
var _S24 : bool;
if(_S23)
{
_S24 = true;
}
else
{
if(k_0 == u32(0))
{
_S24 = (_S18.z) < u32(2);
}
else
{
_S24 = false;
}
}
if(_S24)
{
k_0 = k_0 + u32(1);
continue;
}
var _S25 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S21 + vec3<i32>(vec3<u32>(i_4, j_0, k_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), _S18).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 : vec3<i32> = _S28.id_0 - vec3<i32>(i32(1), 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(1))
{
}
else
{
break;
}
var _S31 : bool;
if(i_5 == u32(0))
{
_S31 = (_S27.x) < u32(2);
}
else
{
_S31 = false;
}
var _S32 : bool;
if(_S31)
{
_S32 = true;
}
else
{
if(j_1 == u32(0))
{
_S32 = (_S27.y) < u32(2);
}
else
{
_S32 = false;
}
}
var _S33 : bool;
if(_S32)
{
_S33 = true;
}
else
{
if(k_1 == u32(0))
{
_S33 = (_S27.z) < u32(2);
}
else
{
_S33 = false;
}
}
if(_S33)
{
k_1 = k_1 + u32(1);
continue;
}
var _S34 : vec3<u32> = vec3<u32>(i_5, j_1, k_1);
var _S35 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S30 + vec3<i32>(_S34)));
var _S36 : vec3<u32> = _S34 * vec3<u32>(u32(4)) + _S27;
var _S37 : u32 = flatten_shared_index_0(_S36.x, _S36.y, _S36.z);
if((_S35.id_1) != u32(4294967295))
{
var _S38 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S35), _S27);
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(1))
{
}
else
{
break;
}
var _S40 : bool;
if(i_6 == u32(0))
{
_S40 = (_S39.x) < u32(2);
}
else
{
_S40 = false;
}
var _S41 : bool;
if(_S40)
{
_S41 = true;
}
else
{
if(j_2 == u32(0))
{
_S41 = (_S39.y) < u32(2);
}
else
{
_S41 = false;
}
}
var _S42 : bool;
if(_S41)
{
_S42 = true;
}
else
{
if(k_2 == u32(0))
{
_S42 = (_S39.z) < u32(2);
}
else
{
_S42 = false;
}
}
if(_S42)
{
k_2 = k_2 + u32(1);
continue;
}
var _S43 : vec3<u32> = vec3<u32>(i_6, j_2, k_2) * vec3<u32>(u32(4)) + _S39;
var _S44 : u32 = flatten_shared_index_0(_S43.x, _S43.y, _S43.z);
var _S45 : u32 = shared_nodes_0[_S44].particle_id_0;
if((shared_nodes_0[_S44].particle_id_0) != u32(4294967295))
{
shared_affinities_0[_S44] = entryPointParams_particles_dyn_0[_S45].cdf_0.affinity_0;
shared_normals_0[_S44] = entryPointParams_particles_dyn_0[_S45].cdf_0.normal_0;
shared_pos_0[_S44].pt_0 = entryPointParams_particles_pos_0[_S45].pt_0;
shared_affine_0[_S44] = mat3x3<f32>(entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(0)][i32(0)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(0)][i32(1)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(0)][i32(2)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(1)][i32(0)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(1)][i32(1)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(1)][i32(2)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(2)][i32(0)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(2)][i32(1)], entryPointParams_particles_dyn_0[_S45].affine_0.data_0[i32(2)][i32(2)]);
shared_vel_mass_0[_S44] = vec4<f32>(entryPointParams_particles_dyn_0[_S45].velocity_0, entryPointParams_particles_dyn_0[_S45].mass_0);
shared_nodes_0[_S44].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S45];
}
else
{
shared_affinities_0[_S44] = u32(0);
var _S46 : vec3<f32> = vec3<f32>(0.0f);
shared_normals_0[_S44] = _S46;
shared_pos_0[_S44].pt_0 = _S46;
shared_affine_0[_S44] = mat3x3<f32>(_S46, _S46, _S46);
shared_vel_mass_0[_S44] = vec4<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( _S47 : u32, _S48 : f32, _S49 : u32, _S50 : u32) -> P2GStepResult_0
{
var impulse_2 : vec3<f32>;
var ang_impulse_2 : vec3<f32>;
var new_momentum_velocity_mass_2 : vec4<f32>;
var _S51 : u32;
var _S52 : vec3<u32>;
var _S53 : vec3<f32>;
var _S54 : u32 = flatten_shared_shift_0(u32(2), u32(2), u32(2));
var _S55 : vec4<f32> = vec4<f32>(0.0f);
var _S56 : vec3<f32> = vec3<f32>(0.0f);
for(;;)
{
for(;;)
{
for(;;)
{
var _S57 : u32 = _S47 - _S54;
_S51 = _S57;
var _S58 : u32 = _S57 + u32(86);
var _S59 : Position_0 = shared_pos_0[_S58];
var _S60 : mat3x3<f32> = shared_affine_0[_S58];
var _S61 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S58], _S48);
var _S62 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S61, _S48);
var _S63 : vec3<f32> = shared_vel_mass_0[_S58].xyz;
var _S64 : f32 = shared_vel_mass_0[_S58].w;
const _S65 : vec3<u32> = vec3<u32>(u32(2), u32(2), u32(2));
_S52 = _S65;
var _S66 : vec3<u32> = _S65 - vec3<u32>(u32(2), u32(2), u32(2));
var _S67 : vec3<f32> = _S63 * vec3<f32>(_S64);
var _S68 : vec3<f32> = vec3<f32>(_S66);
var _S69 : vec3<f32> = vec3<f32>(_S48);
_S53 = _S69;
var _S70 : vec3<f32> = _S61 + _S68 * _S69;
var _S71 : f32 = _S62[i32(0)][_S66.x] * _S62[i32(1)][_S66.y] * _S62[i32(2)][_S66.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S58]))
{
if(_S50 != u32(4294967295))
{
var _S72 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S73 : vec3<f32> = _S70 + _S59.pt_0;
var _S74 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S72, _S73);
var _S75 : vec3<f32> = (_S63 - (_S74 + project_velocity_0(_S63 - _S74, shared_normals_0[_S58]))) * vec3<f32>((_S71 * _S64));
var _S76 : vec3<f32> = cross(_S75, entryPointParams_body_impulses_0[_S50].com_0 - _S73);
new_momentum_velocity_mass_2 = _S55;
ang_impulse_2 = _S76;
impulse_2 = _S75;
break;
}
new_momentum_velocity_mass_2 = _S55;
}
else
{
new_momentum_velocity_mass_2 = vec4<f32>((((_S60) * (_S70))) + _S67, _S64) * vec4<f32>(_S71);
}
ang_impulse_2 = _S56;
impulse_2 = _S56;
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S77 : u32 = _S51 + u32(74);
var _S78 : Position_0 = shared_pos_0[_S77];
var _S79 : mat3x3<f32> = shared_affine_0[_S77];
var _S80 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S77], _S48);
var _S81 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S80, _S48);
var _S82 : vec3<f32> = shared_vel_mass_0[_S77].xyz;
var _S83 : f32 = shared_vel_mass_0[_S77].w;
var _S84 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(0), u32(2));
var _S85 : vec3<f32> = _S82 * vec3<f32>(_S83);
var _S86 : vec3<f32> = _S80 + vec3<f32>(_S84) * _S53;
var _S87 : f32 = _S81[i32(0)][_S84.x] * _S81[i32(1)][_S84.y] * _S81[i32(2)][_S84.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S77]))
{
if(_S50 != u32(4294967295))
{
var _S88 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S89 : vec3<f32> = _S86 + _S78.pt_0;
var _S90 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S88, _S89);
var _S91 : vec3<f32> = (_S82 - (_S90 + project_velocity_0(_S82 - _S90, shared_normals_0[_S77]))) * vec3<f32>((_S87 * _S83));
var impulse_3 : vec3<f32> = impulse_2 + _S91;
ang_impulse_2 = ang_impulse_2 + cross(_S91, entryPointParams_body_impulses_0[_S50].com_0 - _S89);
impulse_2 = impulse_3;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S79) * (_S86))) + _S85, _S83) * vec4<f32>(_S87);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S92 : u32 = _S51 + u32(80);
var _S93 : Position_0 = shared_pos_0[_S92];
var _S94 : mat3x3<f32> = shared_affine_0[_S92];
var _S95 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S92], _S48);
var _S96 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S95, _S48);
var _S97 : vec3<f32> = shared_vel_mass_0[_S92].xyz;
var _S98 : f32 = shared_vel_mass_0[_S92].w;
var _S99 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(1), u32(2));
var _S100 : vec3<f32> = _S97 * vec3<f32>(_S98);
var _S101 : vec3<f32> = _S95 + vec3<f32>(_S99) * _S53;
var _S102 : f32 = _S96[i32(0)][_S99.x] * _S96[i32(1)][_S99.y] * _S96[i32(2)][_S99.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S92]))
{
if(_S50 != u32(4294967295))
{
var _S103 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S104 : vec3<f32> = _S101 + _S93.pt_0;
var _S105 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S103, _S104);
var _S106 : vec3<f32> = (_S97 - (_S105 + project_velocity_0(_S97 - _S105, shared_normals_0[_S92]))) * vec3<f32>((_S102 * _S98));
var impulse_4 : vec3<f32> = impulse_2 + _S106;
ang_impulse_2 = ang_impulse_2 + cross(_S106, entryPointParams_body_impulses_0[_S50].com_0 - _S104);
impulse_2 = impulse_4;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S94) * (_S101))) + _S100, _S98) * vec4<f32>(_S102);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S107 : u32 = _S51 + u32(84);
var _S108 : Position_0 = shared_pos_0[_S107];
var _S109 : mat3x3<f32> = shared_affine_0[_S107];
var _S110 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S107], _S48);
var _S111 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S110, _S48);
var _S112 : vec3<f32> = shared_vel_mass_0[_S107].xyz;
var _S113 : f32 = shared_vel_mass_0[_S107].w;
var _S114 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(2), u32(2));
var _S115 : vec3<f32> = _S112 * vec3<f32>(_S113);
var _S116 : vec3<f32> = _S110 + vec3<f32>(_S114) * _S53;
var _S117 : f32 = _S111[i32(0)][_S114.x] * _S111[i32(1)][_S114.y] * _S111[i32(2)][_S114.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S107]))
{
if(_S50 != u32(4294967295))
{
var _S118 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S119 : vec3<f32> = _S116 + _S108.pt_0;
var _S120 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S118, _S119);
var _S121 : vec3<f32> = (_S112 - (_S120 + project_velocity_0(_S112 - _S120, shared_normals_0[_S107]))) * vec3<f32>((_S117 * _S113));
var impulse_5 : vec3<f32> = impulse_2 + _S121;
ang_impulse_2 = ang_impulse_2 + cross(_S121, entryPointParams_body_impulses_0[_S50].com_0 - _S119);
impulse_2 = impulse_5;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S109) * (_S116))) + _S115, _S113) * vec4<f32>(_S117);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S122 : u32 = _S51 + u32(72);
var _S123 : Position_0 = shared_pos_0[_S122];
var _S124 : mat3x3<f32> = shared_affine_0[_S122];
var _S125 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S122], _S48);
var _S126 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S125, _S48);
var _S127 : vec3<f32> = shared_vel_mass_0[_S122].xyz;
var _S128 : f32 = shared_vel_mass_0[_S122].w;
var _S129 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(0), u32(2));
var _S130 : vec3<f32> = _S127 * vec3<f32>(_S128);
var _S131 : vec3<f32> = _S125 + vec3<f32>(_S129) * _S53;
var _S132 : f32 = _S126[i32(0)][_S129.x] * _S126[i32(1)][_S129.y] * _S126[i32(2)][_S129.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S122]))
{
if(_S50 != u32(4294967295))
{
var _S133 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S134 : vec3<f32> = _S131 + _S123.pt_0;
var _S135 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S133, _S134);
var _S136 : vec3<f32> = (_S127 - (_S135 + project_velocity_0(_S127 - _S135, shared_normals_0[_S122]))) * vec3<f32>((_S132 * _S128));
var impulse_6 : vec3<f32> = impulse_2 + _S136;
ang_impulse_2 = ang_impulse_2 + cross(_S136, entryPointParams_body_impulses_0[_S50].com_0 - _S134);
impulse_2 = impulse_6;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S124) * (_S131))) + _S130, _S128) * vec4<f32>(_S132);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S137 : u32 = _S51 + u32(78);
var _S138 : Position_0 = shared_pos_0[_S137];
var _S139 : mat3x3<f32> = shared_affine_0[_S137];
var _S140 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S137], _S48);
var _S141 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S140, _S48);
var _S142 : vec3<f32> = shared_vel_mass_0[_S137].xyz;
var _S143 : f32 = shared_vel_mass_0[_S137].w;
var _S144 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(1), u32(2));
var _S145 : vec3<f32> = _S142 * vec3<f32>(_S143);
var _S146 : vec3<f32> = _S140 + vec3<f32>(_S144) * _S53;
var _S147 : f32 = _S141[i32(0)][_S144.x] * _S141[i32(1)][_S144.y] * _S141[i32(2)][_S144.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S137]))
{
if(_S50 != u32(4294967295))
{
var _S148 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S149 : vec3<f32> = _S146 + _S138.pt_0;
var _S150 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S148, _S149);
var _S151 : vec3<f32> = (_S142 - (_S150 + project_velocity_0(_S142 - _S150, shared_normals_0[_S137]))) * vec3<f32>((_S147 * _S143));
var impulse_7 : vec3<f32> = impulse_2 + _S151;
ang_impulse_2 = ang_impulse_2 + cross(_S151, entryPointParams_body_impulses_0[_S50].com_0 - _S149);
impulse_2 = impulse_7;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S139) * (_S146))) + _S145, _S143) * vec4<f32>(_S147);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S152 : u32 = _S51 + u32(85);
var _S153 : Position_0 = shared_pos_0[_S152];
var _S154 : mat3x3<f32> = shared_affine_0[_S152];
var _S155 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S152], _S48);
var _S156 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S155, _S48);
var _S157 : vec3<f32> = shared_vel_mass_0[_S152].xyz;
var _S158 : f32 = shared_vel_mass_0[_S152].w;
var _S159 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(2), u32(2));
var _S160 : vec3<f32> = _S157 * vec3<f32>(_S158);
var _S161 : vec3<f32> = _S155 + vec3<f32>(_S159) * _S53;
var _S162 : f32 = _S156[i32(0)][_S159.x] * _S156[i32(1)][_S159.y] * _S156[i32(2)][_S159.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S152]))
{
if(_S50 != u32(4294967295))
{
var _S163 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S164 : vec3<f32> = _S161 + _S153.pt_0;
var _S165 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S163, _S164);
var _S166 : vec3<f32> = (_S157 - (_S165 + project_velocity_0(_S157 - _S165, shared_normals_0[_S152]))) * vec3<f32>((_S162 * _S158));
var impulse_8 : vec3<f32> = impulse_2 + _S166;
ang_impulse_2 = ang_impulse_2 + cross(_S166, entryPointParams_body_impulses_0[_S50].com_0 - _S164);
impulse_2 = impulse_8;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S154) * (_S161))) + _S160, _S158) * vec4<f32>(_S162);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S167 : u32 = _S51 + u32(73);
var _S168 : Position_0 = shared_pos_0[_S167];
var _S169 : mat3x3<f32> = shared_affine_0[_S167];
var _S170 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S167], _S48);
var _S171 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S170, _S48);
var _S172 : vec3<f32> = shared_vel_mass_0[_S167].xyz;
var _S173 : f32 = shared_vel_mass_0[_S167].w;
var _S174 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(0), u32(2));
var _S175 : vec3<f32> = _S172 * vec3<f32>(_S173);
var _S176 : vec3<f32> = _S170 + vec3<f32>(_S174) * _S53;
var _S177 : f32 = _S171[i32(0)][_S174.x] * _S171[i32(1)][_S174.y] * _S171[i32(2)][_S174.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S167]))
{
if(_S50 != u32(4294967295))
{
var _S178 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S179 : vec3<f32> = _S176 + _S168.pt_0;
var _S180 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S178, _S179);
var _S181 : vec3<f32> = (_S172 - (_S180 + project_velocity_0(_S172 - _S180, shared_normals_0[_S167]))) * vec3<f32>((_S177 * _S173));
var impulse_9 : vec3<f32> = impulse_2 + _S181;
ang_impulse_2 = ang_impulse_2 + cross(_S181, entryPointParams_body_impulses_0[_S50].com_0 - _S179);
impulse_2 = impulse_9;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S169) * (_S176))) + _S175, _S173) * vec4<f32>(_S177);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S182 : u32 = _S51 + u32(79);
var _S183 : Position_0 = shared_pos_0[_S182];
var _S184 : mat3x3<f32> = shared_affine_0[_S182];
var _S185 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S182], _S48);
var _S186 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S185, _S48);
var _S187 : vec3<f32> = shared_vel_mass_0[_S182].xyz;
var _S188 : f32 = shared_vel_mass_0[_S182].w;
var _S189 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(1), u32(2));
var _S190 : vec3<f32> = _S187 * vec3<f32>(_S188);
var _S191 : vec3<f32> = _S185 + vec3<f32>(_S189) * _S53;
var _S192 : f32 = _S186[i32(0)][_S189.x] * _S186[i32(1)][_S189.y] * _S186[i32(2)][_S189.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S182]))
{
if(_S50 != u32(4294967295))
{
var _S193 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S194 : vec3<f32> = _S191 + _S183.pt_0;
var _S195 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S193, _S194);
var _S196 : vec3<f32> = (_S187 - (_S195 + project_velocity_0(_S187 - _S195, shared_normals_0[_S182]))) * vec3<f32>((_S192 * _S188));
var impulse_10 : vec3<f32> = impulse_2 + _S196;
ang_impulse_2 = ang_impulse_2 + cross(_S196, entryPointParams_body_impulses_0[_S50].com_0 - _S194);
impulse_2 = impulse_10;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S184) * (_S191))) + _S190, _S188) * vec4<f32>(_S192);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S197 : u32 = _S51 + u32(14);
var _S198 : Position_0 = shared_pos_0[_S197];
var _S199 : mat3x3<f32> = shared_affine_0[_S197];
var _S200 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S197], _S48);
var _S201 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S200, _S48);
var _S202 : vec3<f32> = shared_vel_mass_0[_S197].xyz;
var _S203 : f32 = shared_vel_mass_0[_S197].w;
var _S204 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(2), u32(0));
var _S205 : vec3<f32> = _S202 * vec3<f32>(_S203);
var _S206 : vec3<f32> = _S200 + vec3<f32>(_S204) * _S53;
var _S207 : f32 = _S201[i32(0)][_S204.x] * _S201[i32(1)][_S204.y] * _S201[i32(2)][_S204.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S197]))
{
if(_S50 != u32(4294967295))
{
var _S208 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S209 : vec3<f32> = _S206 + _S198.pt_0;
var _S210 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S208, _S209);
var _S211 : vec3<f32> = (_S202 - (_S210 + project_velocity_0(_S202 - _S210, shared_normals_0[_S197]))) * vec3<f32>((_S207 * _S203));
var impulse_11 : vec3<f32> = impulse_2 + _S211;
ang_impulse_2 = ang_impulse_2 + cross(_S211, entryPointParams_body_impulses_0[_S50].com_0 - _S209);
impulse_2 = impulse_11;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S199) * (_S206))) + _S205, _S203) * vec4<f32>(_S207);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S212 : u32 = _S51 + u32(2);
var _S213 : Position_0 = shared_pos_0[_S212];
var _S214 : mat3x3<f32> = shared_affine_0[_S212];
var _S215 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S212], _S48);
var _S216 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S215, _S48);
var _S217 : vec3<f32> = shared_vel_mass_0[_S212].xyz;
var _S218 : f32 = shared_vel_mass_0[_S212].w;
var _S219 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(0), u32(0));
var _S220 : vec3<f32> = _S217 * vec3<f32>(_S218);
var _S221 : vec3<f32> = _S215 + vec3<f32>(_S219) * _S53;
var _S222 : f32 = _S216[i32(0)][_S219.x] * _S216[i32(1)][_S219.y] * _S216[i32(2)][_S219.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S212]))
{
if(_S50 != u32(4294967295))
{
var _S223 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S224 : vec3<f32> = _S221 + _S213.pt_0;
var _S225 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S223, _S224);
var _S226 : vec3<f32> = (_S217 - (_S225 + project_velocity_0(_S217 - _S225, shared_normals_0[_S212]))) * vec3<f32>((_S222 * _S218));
var impulse_12 : vec3<f32> = impulse_2 + _S226;
ang_impulse_2 = ang_impulse_2 + cross(_S226, entryPointParams_body_impulses_0[_S50].com_0 - _S224);
impulse_2 = impulse_12;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S214) * (_S221))) + _S220, _S218) * vec4<f32>(_S222);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S227 : u32 = _S51 + u32(8);
var _S228 : Position_0 = shared_pos_0[_S227];
var _S229 : mat3x3<f32> = shared_affine_0[_S227];
var _S230 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S227], _S48);
var _S231 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S230, _S48);
var _S232 : vec3<f32> = shared_vel_mass_0[_S227].xyz;
var _S233 : f32 = shared_vel_mass_0[_S227].w;
var _S234 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(1), u32(0));
var _S235 : vec3<f32> = _S232 * vec3<f32>(_S233);
var _S236 : vec3<f32> = _S230 + vec3<f32>(_S234) * _S53;
var _S237 : f32 = _S231[i32(0)][_S234.x] * _S231[i32(1)][_S234.y] * _S231[i32(2)][_S234.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S227]))
{
if(_S50 != u32(4294967295))
{
var _S238 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S239 : vec3<f32> = _S236 + _S228.pt_0;
var _S240 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S238, _S239);
var _S241 : vec3<f32> = (_S232 - (_S240 + project_velocity_0(_S232 - _S240, shared_normals_0[_S227]))) * vec3<f32>((_S237 * _S233));
var impulse_13 : vec3<f32> = impulse_2 + _S241;
ang_impulse_2 = ang_impulse_2 + cross(_S241, entryPointParams_body_impulses_0[_S50].com_0 - _S239);
impulse_2 = impulse_13;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S229) * (_S236))) + _S235, _S233) * vec4<f32>(_S237);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S242 : u32 = _S51 + u32(12);
var _S243 : Position_0 = shared_pos_0[_S242];
var _S244 : mat3x3<f32> = shared_affine_0[_S242];
var _S245 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S242], _S48);
var _S246 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S245, _S48);
var _S247 : vec3<f32> = shared_vel_mass_0[_S242].xyz;
var _S248 : f32 = shared_vel_mass_0[_S242].w;
var _S249 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(2), u32(0));
var _S250 : vec3<f32> = _S247 * vec3<f32>(_S248);
var _S251 : vec3<f32> = _S245 + vec3<f32>(_S249) * _S53;
var _S252 : f32 = _S246[i32(0)][_S249.x] * _S246[i32(1)][_S249.y] * _S246[i32(2)][_S249.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S242]))
{
if(_S50 != u32(4294967295))
{
var _S253 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S254 : vec3<f32> = _S251 + _S243.pt_0;
var _S255 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S253, _S254);
var _S256 : vec3<f32> = (_S247 - (_S255 + project_velocity_0(_S247 - _S255, shared_normals_0[_S242]))) * vec3<f32>((_S252 * _S248));
var impulse_14 : vec3<f32> = impulse_2 + _S256;
ang_impulse_2 = ang_impulse_2 + cross(_S256, entryPointParams_body_impulses_0[_S50].com_0 - _S254);
impulse_2 = impulse_14;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S244) * (_S251))) + _S250, _S248) * vec4<f32>(_S252);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S257 : Position_0 = shared_pos_0[_S51];
var _S258 : mat3x3<f32> = shared_affine_0[_S51];
var _S259 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S51], _S48);
var _S260 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S259, _S48);
var _S261 : vec3<f32> = shared_vel_mass_0[_S51].xyz;
var _S262 : f32 = shared_vel_mass_0[_S51].w;
var _S263 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(0), u32(0));
var _S264 : vec3<f32> = _S261 * vec3<f32>(_S262);
var _S265 : vec3<f32> = _S259 + vec3<f32>(_S263) * _S53;
var _S266 : f32 = _S260[i32(0)][_S263.x] * _S260[i32(1)][_S263.y] * _S260[i32(2)][_S263.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S51]))
{
if(_S50 != u32(4294967295))
{
var _S267 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S268 : vec3<f32> = _S265 + _S257.pt_0;
var _S269 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S267, _S268);
var _S270 : vec3<f32> = (_S261 - (_S269 + project_velocity_0(_S261 - _S269, shared_normals_0[_S51]))) * vec3<f32>((_S266 * _S262));
var impulse_15 : vec3<f32> = impulse_2 + _S270;
ang_impulse_2 = ang_impulse_2 + cross(_S270, entryPointParams_body_impulses_0[_S50].com_0 - _S268);
impulse_2 = impulse_15;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S258) * (_S265))) + _S264, _S262) * vec4<f32>(_S266);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S271 : u32 = _S51 + u32(6);
var _S272 : Position_0 = shared_pos_0[_S271];
var _S273 : mat3x3<f32> = shared_affine_0[_S271];
var _S274 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S271], _S48);
var _S275 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S274, _S48);
var _S276 : vec3<f32> = shared_vel_mass_0[_S271].xyz;
var _S277 : f32 = shared_vel_mass_0[_S271].w;
var _S278 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(1), u32(0));
var _S279 : vec3<f32> = _S276 * vec3<f32>(_S277);
var _S280 : vec3<f32> = _S274 + vec3<f32>(_S278) * _S53;
var _S281 : f32 = _S275[i32(0)][_S278.x] * _S275[i32(1)][_S278.y] * _S275[i32(2)][_S278.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S271]))
{
if(_S50 != u32(4294967295))
{
var _S282 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S283 : vec3<f32> = _S280 + _S272.pt_0;
var _S284 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S282, _S283);
var _S285 : vec3<f32> = (_S276 - (_S284 + project_velocity_0(_S276 - _S284, shared_normals_0[_S271]))) * vec3<f32>((_S281 * _S277));
var impulse_16 : vec3<f32> = impulse_2 + _S285;
ang_impulse_2 = ang_impulse_2 + cross(_S285, entryPointParams_body_impulses_0[_S50].com_0 - _S283);
impulse_2 = impulse_16;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S273) * (_S280))) + _S279, _S277) * vec4<f32>(_S281);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S286 : u32 = _S51 + u32(13);
var _S287 : Position_0 = shared_pos_0[_S286];
var _S288 : mat3x3<f32> = shared_affine_0[_S286];
var _S289 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S286], _S48);
var _S290 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S289, _S48);
var _S291 : vec3<f32> = shared_vel_mass_0[_S286].xyz;
var _S292 : f32 = shared_vel_mass_0[_S286].w;
var _S293 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(2), u32(0));
var _S294 : vec3<f32> = _S291 * vec3<f32>(_S292);
var _S295 : vec3<f32> = _S289 + vec3<f32>(_S293) * _S53;
var _S296 : f32 = _S290[i32(0)][_S293.x] * _S290[i32(1)][_S293.y] * _S290[i32(2)][_S293.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S286]))
{
if(_S50 != u32(4294967295))
{
var _S297 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S298 : vec3<f32> = _S295 + _S287.pt_0;
var _S299 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S297, _S298);
var _S300 : vec3<f32> = (_S291 - (_S299 + project_velocity_0(_S291 - _S299, shared_normals_0[_S286]))) * vec3<f32>((_S296 * _S292));
var impulse_17 : vec3<f32> = impulse_2 + _S300;
ang_impulse_2 = ang_impulse_2 + cross(_S300, entryPointParams_body_impulses_0[_S50].com_0 - _S298);
impulse_2 = impulse_17;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S288) * (_S295))) + _S294, _S292) * vec4<f32>(_S296);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S301 : u32 = _S51 + u32(1);
var _S302 : Position_0 = shared_pos_0[_S301];
var _S303 : mat3x3<f32> = shared_affine_0[_S301];
var _S304 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S301], _S48);
var _S305 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S304, _S48);
var _S306 : vec3<f32> = shared_vel_mass_0[_S301].xyz;
var _S307 : f32 = shared_vel_mass_0[_S301].w;
var _S308 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(0), u32(0));
var _S309 : vec3<f32> = _S306 * vec3<f32>(_S307);
var _S310 : vec3<f32> = _S304 + vec3<f32>(_S308) * _S53;
var _S311 : f32 = _S305[i32(0)][_S308.x] * _S305[i32(1)][_S308.y] * _S305[i32(2)][_S308.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S301]))
{
if(_S50 != u32(4294967295))
{
var _S312 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S313 : vec3<f32> = _S310 + _S302.pt_0;
var _S314 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S312, _S313);
var _S315 : vec3<f32> = (_S306 - (_S314 + project_velocity_0(_S306 - _S314, shared_normals_0[_S301]))) * vec3<f32>((_S311 * _S307));
var impulse_18 : vec3<f32> = impulse_2 + _S315;
ang_impulse_2 = ang_impulse_2 + cross(_S315, entryPointParams_body_impulses_0[_S50].com_0 - _S313);
impulse_2 = impulse_18;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S303) * (_S310))) + _S309, _S307) * vec4<f32>(_S311);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S316 : u32 = _S51 + u32(7);
var _S317 : Position_0 = shared_pos_0[_S316];
var _S318 : mat3x3<f32> = shared_affine_0[_S316];
var _S319 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S316], _S48);
var _S320 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S319, _S48);
var _S321 : vec3<f32> = shared_vel_mass_0[_S316].xyz;
var _S322 : f32 = shared_vel_mass_0[_S316].w;
var _S323 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(1), u32(0));
var _S324 : vec3<f32> = _S321 * vec3<f32>(_S322);
var _S325 : vec3<f32> = _S319 + vec3<f32>(_S323) * _S53;
var _S326 : f32 = _S320[i32(0)][_S323.x] * _S320[i32(1)][_S323.y] * _S320[i32(2)][_S323.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S316]))
{
if(_S50 != u32(4294967295))
{
var _S327 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S328 : vec3<f32> = _S325 + _S317.pt_0;
var _S329 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S327, _S328);
var _S330 : vec3<f32> = (_S321 - (_S329 + project_velocity_0(_S321 - _S329, shared_normals_0[_S316]))) * vec3<f32>((_S326 * _S322));
var impulse_19 : vec3<f32> = impulse_2 + _S330;
ang_impulse_2 = ang_impulse_2 + cross(_S330, entryPointParams_body_impulses_0[_S50].com_0 - _S328);
impulse_2 = impulse_19;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S318) * (_S325))) + _S324, _S322) * vec4<f32>(_S326);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S331 : u32 = _S51 + u32(50);
var _S332 : Position_0 = shared_pos_0[_S331];
var _S333 : mat3x3<f32> = shared_affine_0[_S331];
var _S334 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S331], _S48);
var _S335 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S334, _S48);
var _S336 : vec3<f32> = shared_vel_mass_0[_S331].xyz;
var _S337 : f32 = shared_vel_mass_0[_S331].w;
var _S338 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(2), u32(1));
var _S339 : vec3<f32> = _S336 * vec3<f32>(_S337);
var _S340 : vec3<f32> = _S334 + vec3<f32>(_S338) * _S53;
var _S341 : f32 = _S335[i32(0)][_S338.x] * _S335[i32(1)][_S338.y] * _S335[i32(2)][_S338.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S331]))
{
if(_S50 != u32(4294967295))
{
var _S342 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S343 : vec3<f32> = _S340 + _S332.pt_0;
var _S344 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S342, _S343);
var _S345 : vec3<f32> = (_S336 - (_S344 + project_velocity_0(_S336 - _S344, shared_normals_0[_S331]))) * vec3<f32>((_S341 * _S337));
var impulse_20 : vec3<f32> = impulse_2 + _S345;
ang_impulse_2 = ang_impulse_2 + cross(_S345, entryPointParams_body_impulses_0[_S50].com_0 - _S343);
impulse_2 = impulse_20;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S333) * (_S340))) + _S339, _S337) * vec4<f32>(_S341);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S346 : u32 = _S51 + u32(38);
var _S347 : Position_0 = shared_pos_0[_S346];
var _S348 : mat3x3<f32> = shared_affine_0[_S346];
var _S349 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S346], _S48);
var _S350 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S349, _S48);
var _S351 : vec3<f32> = shared_vel_mass_0[_S346].xyz;
var _S352 : f32 = shared_vel_mass_0[_S346].w;
var _S353 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(0), u32(1));
var _S354 : vec3<f32> = _S351 * vec3<f32>(_S352);
var _S355 : vec3<f32> = _S349 + vec3<f32>(_S353) * _S53;
var _S356 : f32 = _S350[i32(0)][_S353.x] * _S350[i32(1)][_S353.y] * _S350[i32(2)][_S353.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S346]))
{
if(_S50 != u32(4294967295))
{
var _S357 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S358 : vec3<f32> = _S355 + _S347.pt_0;
var _S359 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S357, _S358);
var _S360 : vec3<f32> = (_S351 - (_S359 + project_velocity_0(_S351 - _S359, shared_normals_0[_S346]))) * vec3<f32>((_S356 * _S352));
var impulse_21 : vec3<f32> = impulse_2 + _S360;
ang_impulse_2 = ang_impulse_2 + cross(_S360, entryPointParams_body_impulses_0[_S50].com_0 - _S358);
impulse_2 = impulse_21;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S348) * (_S355))) + _S354, _S352) * vec4<f32>(_S356);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S361 : u32 = _S51 + u32(44);
var _S362 : Position_0 = shared_pos_0[_S361];
var _S363 : mat3x3<f32> = shared_affine_0[_S361];
var _S364 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S361], _S48);
var _S365 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S364, _S48);
var _S366 : vec3<f32> = shared_vel_mass_0[_S361].xyz;
var _S367 : f32 = shared_vel_mass_0[_S361].w;
var _S368 : vec3<u32> = _S52 - vec3<u32>(u32(2), u32(1), u32(1));
var _S369 : vec3<f32> = _S366 * vec3<f32>(_S367);
var _S370 : vec3<f32> = _S364 + vec3<f32>(_S368) * _S53;
var _S371 : f32 = _S365[i32(0)][_S368.x] * _S365[i32(1)][_S368.y] * _S365[i32(2)][_S368.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S361]))
{
if(_S50 != u32(4294967295))
{
var _S372 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S373 : vec3<f32> = _S370 + _S362.pt_0;
var _S374 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S372, _S373);
var _S375 : vec3<f32> = (_S366 - (_S374 + project_velocity_0(_S366 - _S374, shared_normals_0[_S361]))) * vec3<f32>((_S371 * _S367));
var impulse_22 : vec3<f32> = impulse_2 + _S375;
ang_impulse_2 = ang_impulse_2 + cross(_S375, entryPointParams_body_impulses_0[_S50].com_0 - _S373);
impulse_2 = impulse_22;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S363) * (_S370))) + _S369, _S367) * vec4<f32>(_S371);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S376 : u32 = _S51 + u32(48);
var _S377 : Position_0 = shared_pos_0[_S376];
var _S378 : mat3x3<f32> = shared_affine_0[_S376];
var _S379 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S376], _S48);
var _S380 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S379, _S48);
var _S381 : vec3<f32> = shared_vel_mass_0[_S376].xyz;
var _S382 : f32 = shared_vel_mass_0[_S376].w;
var _S383 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(2), u32(1));
var _S384 : vec3<f32> = _S381 * vec3<f32>(_S382);
var _S385 : vec3<f32> = _S379 + vec3<f32>(_S383) * _S53;
var _S386 : f32 = _S380[i32(0)][_S383.x] * _S380[i32(1)][_S383.y] * _S380[i32(2)][_S383.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S376]))
{
if(_S50 != u32(4294967295))
{
var _S387 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S388 : vec3<f32> = _S385 + _S377.pt_0;
var _S389 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S387, _S388);
var _S390 : vec3<f32> = (_S381 - (_S389 + project_velocity_0(_S381 - _S389, shared_normals_0[_S376]))) * vec3<f32>((_S386 * _S382));
var impulse_23 : vec3<f32> = impulse_2 + _S390;
ang_impulse_2 = ang_impulse_2 + cross(_S390, entryPointParams_body_impulses_0[_S50].com_0 - _S388);
impulse_2 = impulse_23;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S378) * (_S385))) + _S384, _S382) * vec4<f32>(_S386);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S391 : u32 = _S51 + u32(36);
var _S392 : Position_0 = shared_pos_0[_S391];
var _S393 : mat3x3<f32> = shared_affine_0[_S391];
var _S394 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S391], _S48);
var _S395 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S394, _S48);
var _S396 : vec3<f32> = shared_vel_mass_0[_S391].xyz;
var _S397 : f32 = shared_vel_mass_0[_S391].w;
var _S398 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(0), u32(1));
var _S399 : vec3<f32> = _S396 * vec3<f32>(_S397);
var _S400 : vec3<f32> = _S394 + vec3<f32>(_S398) * _S53;
var _S401 : f32 = _S395[i32(0)][_S398.x] * _S395[i32(1)][_S398.y] * _S395[i32(2)][_S398.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S391]))
{
if(_S50 != u32(4294967295))
{
var _S402 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S403 : vec3<f32> = _S400 + _S392.pt_0;
var _S404 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S402, _S403);
var _S405 : vec3<f32> = (_S396 - (_S404 + project_velocity_0(_S396 - _S404, shared_normals_0[_S391]))) * vec3<f32>((_S401 * _S397));
var impulse_24 : vec3<f32> = impulse_2 + _S405;
ang_impulse_2 = ang_impulse_2 + cross(_S405, entryPointParams_body_impulses_0[_S50].com_0 - _S403);
impulse_2 = impulse_24;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S393) * (_S400))) + _S399, _S397) * vec4<f32>(_S401);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S406 : u32 = _S51 + u32(42);
var _S407 : Position_0 = shared_pos_0[_S406];
var _S408 : mat3x3<f32> = shared_affine_0[_S406];
var _S409 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S406], _S48);
var _S410 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S409, _S48);
var _S411 : vec3<f32> = shared_vel_mass_0[_S406].xyz;
var _S412 : f32 = shared_vel_mass_0[_S406].w;
var _S413 : vec3<u32> = _S52 - vec3<u32>(u32(0), u32(1), u32(1));
var _S414 : vec3<f32> = _S411 * vec3<f32>(_S412);
var _S415 : vec3<f32> = _S409 + vec3<f32>(_S413) * _S53;
var _S416 : f32 = _S410[i32(0)][_S413.x] * _S410[i32(1)][_S413.y] * _S410[i32(2)][_S413.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S406]))
{
if(_S50 != u32(4294967295))
{
var _S417 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S418 : vec3<f32> = _S415 + _S407.pt_0;
var _S419 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S417, _S418);
var _S420 : vec3<f32> = (_S411 - (_S419 + project_velocity_0(_S411 - _S419, shared_normals_0[_S406]))) * vec3<f32>((_S416 * _S412));
var impulse_25 : vec3<f32> = impulse_2 + _S420;
ang_impulse_2 = ang_impulse_2 + cross(_S420, entryPointParams_body_impulses_0[_S50].com_0 - _S418);
impulse_2 = impulse_25;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S408) * (_S415))) + _S414, _S412) * vec4<f32>(_S416);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S421 : u32 = _S51 + u32(49);
var _S422 : Position_0 = shared_pos_0[_S421];
var _S423 : mat3x3<f32> = shared_affine_0[_S421];
var _S424 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S421], _S48);
var _S425 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S424, _S48);
var _S426 : vec3<f32> = shared_vel_mass_0[_S421].xyz;
var _S427 : f32 = shared_vel_mass_0[_S421].w;
var _S428 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(2), u32(1));
var _S429 : vec3<f32> = _S426 * vec3<f32>(_S427);
var _S430 : vec3<f32> = _S424 + vec3<f32>(_S428) * _S53;
var _S431 : f32 = _S425[i32(0)][_S428.x] * _S425[i32(1)][_S428.y] * _S425[i32(2)][_S428.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S421]))
{
if(_S50 != u32(4294967295))
{
var _S432 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S433 : vec3<f32> = _S430 + _S422.pt_0;
var _S434 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S432, _S433);
var _S435 : vec3<f32> = (_S426 - (_S434 + project_velocity_0(_S426 - _S434, shared_normals_0[_S421]))) * vec3<f32>((_S431 * _S427));
var impulse_26 : vec3<f32> = impulse_2 + _S435;
ang_impulse_2 = ang_impulse_2 + cross(_S435, entryPointParams_body_impulses_0[_S50].com_0 - _S433);
impulse_2 = impulse_26;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S423) * (_S430))) + _S429, _S427) * vec4<f32>(_S431);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S436 : u32 = _S51 + u32(37);
var _S437 : Position_0 = shared_pos_0[_S436];
var _S438 : mat3x3<f32> = shared_affine_0[_S436];
var _S439 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S436], _S48);
var _S440 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S439, _S48);
var _S441 : vec3<f32> = shared_vel_mass_0[_S436].xyz;
var _S442 : f32 = shared_vel_mass_0[_S436].w;
var _S443 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(0), u32(1));
var _S444 : vec3<f32> = _S441 * vec3<f32>(_S442);
var _S445 : vec3<f32> = _S439 + vec3<f32>(_S443) * _S53;
var _S446 : f32 = _S440[i32(0)][_S443.x] * _S440[i32(1)][_S443.y] * _S440[i32(2)][_S443.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S436]))
{
if(_S50 != u32(4294967295))
{
var _S447 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S448 : vec3<f32> = _S445 + _S437.pt_0;
var _S449 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S447, _S448);
var _S450 : vec3<f32> = (_S441 - (_S449 + project_velocity_0(_S441 - _S449, shared_normals_0[_S436]))) * vec3<f32>((_S446 * _S442));
var impulse_27 : vec3<f32> = impulse_2 + _S450;
ang_impulse_2 = ang_impulse_2 + cross(_S450, entryPointParams_body_impulses_0[_S50].com_0 - _S448);
impulse_2 = impulse_27;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S438) * (_S445))) + _S444, _S442) * vec4<f32>(_S446);
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S451 : u32 = _S51 + u32(43);
var _S452 : Position_0 = shared_pos_0[_S451];
var _S453 : mat3x3<f32> = shared_affine_0[_S451];
var _S454 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S451], _S48);
var _S455 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S454, _S48);
var _S456 : vec3<f32> = shared_vel_mass_0[_S451].xyz;
var _S457 : f32 = shared_vel_mass_0[_S451].w;
var _S458 : vec3<u32> = _S52 - vec3<u32>(u32(1), u32(1), u32(1));
var _S459 : vec3<f32> = _S456 * vec3<f32>(_S457);
var _S460 : vec3<f32> = _S454 + vec3<f32>(_S458) * _S53;
var _S461 : f32 = _S455[i32(0)][_S458.x] * _S455[i32(1)][_S458.y] * _S455[i32(2)][_S458.z];
if(!affinities_are_compatible_0(_S49, shared_affinities_0[_S451]))
{
if(_S50 != u32(4294967295))
{
var _S462 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S50].linear_0, entryPointParams_body_vels_0[_S50].angular_0 );
var _S463 : vec3<f32> = _S460 + _S452.pt_0;
var _S464 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S50].com_0, _S462, _S463);
var _S465 : vec3<f32> = (_S456 - (_S464 + project_velocity_0(_S456 - _S464, shared_normals_0[_S451]))) * vec3<f32>((_S461 * _S457));
var impulse_28 : vec3<f32> = impulse_2 + _S465;
ang_impulse_2 = ang_impulse_2 + cross(_S465, entryPointParams_body_impulses_0[_S50].com_0 - _S463);
impulse_2 = impulse_28;
break;
}
}
else
{
new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S453) * (_S460))) + _S459, _S457) * vec4<f32>(_S461);
}
break;
}
break;
}
break;
}
return P2GStepResult_x24init_0(new_momentum_velocity_mass_2, impulse_2, ang_impulse_2);
}
@compute
@workgroup_size(4, 4, 4)
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 _S466 : u32 = block_id_0.x;
var _S467 : BlockVirtualId_0 = BlockVirtualId_0( entryPointParams_active_blocks_0[_S466].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, _S467, _S466);
workgroupBarrier();
var _S468 : u32 = atomicLoad(&(max_linked_list_length_0));
max_linked_list_length_uniform_0 = _S468;
fetch_nodes_0(tid_0, _S467, _S466);
var _S469 : u32 = flatten_shared_index_0(tid_0.x + u32(4), tid_0.y + u32(4), tid_0.z + u32(4));
var _S470 : u32 = shared_nodes_0[_S469].global_id_0;
var _S471 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S469].global_id_0].cdf_1.affinities_0;
var _S472 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S469].global_id_0].cdf_1.closest_id_0;
const _S473 : vec3<f32> = vec3<f32>(0.0f, 0.0f, 0.0f);
var total_result_0 : P2GStepResult_0 = P2GStepResult_x24init_0(vec4<f32>(0.0f, 0.0f, 0.0f, 0.0f), _S473, _S473);
var _S474 : u32 = (workgroupUniformLoad(&((max_linked_list_length_uniform_0))));
var i_7 : u32 = u32(0);
for(;;)
{
if(i_7 < _S474)
{
}
else
{
break;
}
workgroupBarrier();
fetch_next_particle_0(tid_0);
workgroupBarrier();
var _S475 : P2GStepResult_0 = p2g_step_0(_S469, entryPointParams_grid_0[i32(0)].cell_width_0, _S471, _S472);
total_result_0.new_momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0 + _S475.new_momentum_velocity_mass_0;
total_result_0.impulse_0 = total_result_0.impulse_0 + _S475.impulse_0;
total_result_0.ang_impulse_0 = total_result_0.ang_impulse_0 + _S475.ang_impulse_0;
i_7 = i_7 + u32(1);
}
entryPointParams_nodes_0[_S470].momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0;
if(_S472 != u32(4294967295))
{
var _S476 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].linear_x_0), flt2int_0(total_result_0.impulse_0.x));
var _S477 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].linear_y_0), flt2int_0(total_result_0.impulse_0.y));
var _S478 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].linear_z_0), flt2int_0(total_result_0.impulse_0.z));
var _S479 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].angular_x_0), flt2int_0(total_result_0.ang_impulse_0.x));
var _S480 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].angular_y_0), flt2int_0(total_result_0.ang_impulse_0.y));
var _S481 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S472].angular_z_0), flt2int_0(total_result_0.ang_impulse_0.z));
}
return;
}