struct SimulationParams_std140_0
{
@align(16) gravity_0 : vec2<f32>,
@align(8) padding_0 : f32,
@align(4) dt_0 : f32,
};
@binding(0) @group(0) var<uniform> entryPointParams_params_0 : SimulationParams_std140_0;
struct GridGeneric_std430_0
{
@align(4) num_active_blocks_0 : u32,
@align(4) cell_width_0 : f32,
@align(4) hmap_capacity_0 : u32,
@align(4) capacity_0 : u32,
};
@binding(1) @group(0) var<storage, read> entryPointParams_grid_0 : array<GridGeneric_std430_0>;
struct BlockVirtualId_std430_0
{
@align(8) id_0 : vec2<i32>,
};
struct 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(2) @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(3) @group(0) var<storage, read> entryPointParams_active_blocks_0 : array<ActiveBlockHeaderGeneric_std430_0>;
struct NodeCdf_std430_0
{
@align(4) distance_0 : f32,
@align(4) affinities_0 : u32,
@align(4) closest_id_0 : u32,
};
struct Node_std430_0
{
@align(16) momentum_velocity_mass_0 : vec3<f32>,
@align(4) cdf_0 : NodeCdf_std430_0,
};
@binding(4) @group(0) var<storage, read> entryPointParams_nodes_0 : array<Node_std430_0>;
@binding(5) @group(0) var<storage, read> entryPointParams_sorted_particle_ids_0 : array<u32>;
struct Position_std430_0
{
@align(8) pt_0 : vec2<f32>,
};
@binding(6) @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_1 : 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(7) @group(0) var<storage, read_write> entryPointParams_particles_dyn_0 : array<Dynamics_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 BodyMassProperties_std430_0
{
@align(8) inv_inertia_0 : f32,
@align(8) inv_mass_0 : vec2<f32>,
@align(8) com_0 : vec2<f32>,
};
@binding(9) @group(0) var<storage, read> entryPointParams_body_mprops_0 : array<BodyMassProperties_std430_0>;
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;
}
fn flatten_shared_index_0( x_0 : u32, y_0 : u32) -> u32
{
return x_0 + y_0 * u32(10);
}
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> shared_nodes_vel_mass_0 : array<vec3<f32>, i32(100)>;
struct NodeCdf_0
{
distance_0 : f32,
affinities_0 : u32,
closest_id_0 : u32,
};
var<workgroup> shared_nodes_cdf_0 : array<NodeCdf_0, i32(100)>;
fn NodeCdf_x24init_0( distance_1 : f32, affinities_1 : u32, closest_id_1 : u32) -> NodeCdf_0
{
var _S5 : NodeCdf_0;
_S5.distance_0 = distance_1;
_S5.affinities_0 = affinities_1;
_S5.closest_id_0 = closest_id_1;
return _S5;
}
fn QuadraticKernel_inv_d_0( cell_width_1 : f32) -> f32
{
return 4.0f / (cell_width_1 * cell_width_1);
}
fn associated_grid_pos_0( part_pos_0 : ptr<function, Position_std430_0>, cell_width_2 : f32) -> vec2<f32>
{
var _S6 : vec2<f32> = vec2<f32>(cell_width_2);
return (round((*part_pos_0).pt_0 / _S6) - vec2<f32>(1.0f)) * _S6;
}
fn dir_to_associated_grid_node_0( part_pos_1 : ptr<function, Position_std430_0>, cell_width_3 : f32) -> vec2<f32>
{
var _S7 : vec2<f32> = associated_grid_pos_0(&((*part_pos_1)), cell_width_3);
return _S7 - (*part_pos_1).pt_0;
}
fn QuadraticKernel_eval_all_0( x_1 : f32) -> vec3<f32>
{
var _S8 : f32 = 1.5f - x_1;
var _S9 : f32 = x_1 - 1.0f;
var _S10 : f32 = x_1 - 0.5f;
return vec3<f32>(0.5f * _S8 * _S8, 0.75f - _S9 * _S9, 0.5f * _S10 * _S10);
}
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 associated_cell_index_in_block_off_by_one_0( part_pos_2 : ptr<function, Position_std430_0>, cell_width_4 : f32) -> vec2<u32>
{
var _S11 : vec2<f32> = round((*part_pos_2).pt_0 / vec2<f32>(cell_width_4)) - vec2<f32>(1.0f);
var _S12 : vec2<f32> = vec2<f32>(8.0f);
return vec2<u32>(_S11 - floor(_S11 / _S12) * _S12);
}
fn affinities_are_compatible_0( affinity1_0 : u32, affinity2_0 : u32) -> bool
{
var _S13 : u32 = (((affinity1_0 & (affinity2_0))) & (u32(65535)));
return ((((affinity1_0 >> (u32(16)))) & (_S13))) == ((((affinity2_0 >> (u32(16)))) & (_S13)));
}
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 _S14 : vec2<f32> = point_0 - center_of_mass_0;
return (*vels_0).linear_0 + vec2<f32>((*vels_0).angular_0) * vec2<f32>(- _S14.y, _S14.x);
}
fn project_velocity_0( vel_0 : vec2<f32>, n_0 : vec2<f32>) -> vec2<f32>
{
var _S15 : f32 = dot(vel_0, n_0);
if(_S15 < 0.0f)
{
var _S16 : vec2<f32> = vel_0 - n_0 * vec2<f32>(_S15);
var _S17 : f32 = length(_S16);
return select(vec2<f32>(0.0f), _S16 / vec2<f32>(_S17), _S17 > 9.99999993922529029e-09f) * vec2<f32>(max(0.0f, _S17 + 20.0f * _S15));
}
else
{
return vel_0;
}
}
fn outer_product_0( a_0 : vec2<f32>, b_0 : vec2<f32>) -> mat2x2<f32>
{
return mat2x2<f32>(a_0 * vec2<f32>(b_0.x), a_0 * vec2<f32>(b_0.y));
}
fn affinity_bit_0( i_collider_0 : u32, affinity_1 : u32) -> bool
{
return ((affinity_1 & (((u32(1) << (i_collider_0)))))) != u32(0);
}
fn find_block_header_id_0( _S18 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S19 : u32 = pack_key_0(_S18);
var slot_0 : u32 = ((hash_0(_S19)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S20 : GridHashMapEntryGeneric_std430_0 = entryPointParams_hmap_entries_0[slot_0];
if((_S20.state_0) == _S19)
{
var _S21 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S21;
}
else
{
if((_S20.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 global_shared_memory_transfers_0( _S22 : vec3<u32>, _S23 : ptr<function, BlockVirtualId_std430_0>)
{
var _S24 : vec2<i32> = (*_S23).id_0;
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 _S25 : bool;
if(i_4 == u32(1))
{
_S25 = (_S22.x) > u32(1);
}
else
{
_S25 = false;
}
var _S26 : bool;
if(_S25)
{
_S26 = true;
}
else
{
if(j_0 == u32(1))
{
_S26 = (_S22.y) > u32(1);
}
else
{
_S26 = false;
}
}
if(_S26)
{
j_0 = j_0 + u32(1);
continue;
}
var _S27 : vec2<u32> = vec2<u32>(i_4, j_0);
var _S28 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S24 + vec2<i32>(_S27)));
var _S29 : vec2<u32> = _S22.xy;
var _S30 : vec2<u32> = _S27 * vec2<u32>(u32(8)) + _S29;
var _S31 : u32 = flatten_shared_index_0(_S30.x, _S30.y);
if((_S28.id_1) != u32(4294967295))
{
var _S32 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S28), _S29);
shared_nodes_vel_mass_0[_S31] = entryPointParams_nodes_0[_S32.id_3].momentum_velocity_mass_0;
var _S33 : Node_std430_0 = entryPointParams_nodes_0[_S32.id_3];
shared_nodes_cdf_0[_S31].distance_0 = _S33.cdf_0.distance_0;
shared_nodes_cdf_0[_S31].affinities_0 = _S33.cdf_0.affinities_0;
shared_nodes_cdf_0[_S31].closest_id_0 = _S33.cdf_0.closest_id_0;
}
else
{
shared_nodes_vel_mass_0[_S31] = vec3<f32>(0.0f);
shared_nodes_cdf_0[_S31] = NodeCdf_x24init_0(0.0f, u32(0), u32(4294967295));
}
j_0 = j_0 + u32(1);
}
i_4 = i_4 + u32(1);
}
return;
}
fn particle_g2p_0( _S34 : u32, _S35 : f32, _S36 : f32)
{
var _S37 : vec2<f32> = vec2<f32>(0.0f);
var _S38 : vec3<f32> = vec3<f32>(0.0f);
var _S39 : mat2x2<f32> = mat2x2<f32>(_S37, _S37);
var momentum_velocity_mass_1 : vec3<f32>;
var rigid_vel_1 : vec2<f32>;
var velocity_gradient_0 : mat2x2<f32>;
var vel_grad_det_1 : f32;
if((entryPointParams_particles_dyn_0[_S34].enabled_0) != u32(0))
{
var _S40 : Position_std430_0 = entryPointParams_particles_pos_0[_S34];
var _S41 : vec2<f32> = entryPointParams_particles_dyn_0[_S34].velocity_0;
var _S42 : Cdf_std430_0 = entryPointParams_particles_dyn_0[_S34].cdf_1;
var _S43 : f32 = QuadraticKernel_inv_d_0(_S35);
var _S44 : vec2<f32> = dir_to_associated_grid_node_0(&(_S40), _S35);
var _S45 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S44, _S35);
var _S46 : vec2<u32> = associated_cell_index_in_block_off_by_one_0(&(_S40), _S35);
var _S47 : u32 = flatten_shared_index_0(_S46.x, _S46.y);
var _S48 : u32 = _S47 + u32(22);
var _S49 : vec3<f32> = shared_nodes_vel_mass_0[_S48];
var _S50 : NodeCdf_0 = shared_nodes_cdf_0[_S48];
var _S51 : vec2<f32> = vec2<f32>(_S35);
var _S52 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(2), u32(2))) * _S51;
if(!affinities_are_compatible_0(entryPointParams_particles_dyn_0[_S34].cdf_1.affinity_0, shared_nodes_cdf_0[_S48].affinities_0))
{
if((_S50.closest_id_0) != u32(4294967295))
{
var _S53 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S50.closest_id_0];
var _S54 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S50.closest_id_0].com_0, &(_S53), _S52 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S54 + project_velocity_0(_S41 - _S54, _S42.normal_0), _S49.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S49.z);
}
}
else
{
momentum_velocity_mass_1 = _S49;
}
var _S55 : f32 = _S45[i32(0)][u32(2)] * _S45[i32(1)][u32(2)];
var _S56 : vec3<f32> = momentum_velocity_mass_1 * vec3<f32>(_S55);
var _S57 : f32 = _S55 * _S43;
var _S58 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S59 : mat2x2<f32> = outer_product_0(_S58, _S52);
var _S60 : mat2x2<f32> = mat2x2<f32>(_S57, _S57, _S57, _S57);
var _S61 : mat2x2<f32> = mat2x2<f32>(_S60[0] * _S59[0], _S60[1] * _S59[1]);
var _S62 : f32 = _S57 * dot(_S58, _S52);
var _S63 : u32 = _S47 + u32(2);
var _S64 : vec3<f32> = shared_nodes_vel_mass_0[_S63];
var _S65 : NodeCdf_0 = shared_nodes_cdf_0[_S63];
var _S66 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(2), u32(0))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S63].affinities_0))
{
if((_S65.closest_id_0) != u32(4294967295))
{
var _S67 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S65.closest_id_0];
var _S68 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S65.closest_id_0].com_0, &(_S67), _S66 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S68 + project_velocity_0(_S41 - _S68, _S42.normal_0), _S64.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S64.z);
}
}
else
{
momentum_velocity_mass_1 = _S64;
}
var _S69 : f32 = _S45[i32(0)][u32(2)] * _S45[i32(1)][u32(0)];
var momentum_velocity_mass_2 : vec3<f32> = _S56 + momentum_velocity_mass_1 * vec3<f32>(_S69);
var _S70 : f32 = _S69 * _S43;
var _S71 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S72 : mat2x2<f32> = outer_product_0(_S71, _S66);
var _S73 : mat2x2<f32> = mat2x2<f32>(_S70, _S70, _S70, _S70);
var _S74 : mat2x2<f32> = _S61 + mat2x2<f32>(_S73[0] * _S72[0], _S73[1] * _S72[1]);
var vel_grad_det_2 : f32 = _S62 + _S70 * dot(_S71, _S66);
var _S75 : u32 = _S47 + u32(12);
var _S76 : vec3<f32> = shared_nodes_vel_mass_0[_S75];
var _S77 : NodeCdf_0 = shared_nodes_cdf_0[_S75];
var _S78 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(2), u32(1))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S75].affinities_0))
{
if((_S77.closest_id_0) != u32(4294967295))
{
var _S79 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S77.closest_id_0];
var _S80 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S77.closest_id_0].com_0, &(_S79), _S78 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S80 + project_velocity_0(_S41 - _S80, _S42.normal_0), _S76.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S76.z);
}
}
else
{
momentum_velocity_mass_1 = _S76;
}
var _S81 : f32 = _S45[i32(0)][u32(2)] * _S45[i32(1)][u32(1)];
var momentum_velocity_mass_3 : vec3<f32> = momentum_velocity_mass_2 + momentum_velocity_mass_1 * vec3<f32>(_S81);
var _S82 : f32 = _S81 * _S43;
var _S83 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S84 : mat2x2<f32> = outer_product_0(_S83, _S78);
var _S85 : mat2x2<f32> = mat2x2<f32>(_S82, _S82, _S82, _S82);
var _S86 : mat2x2<f32> = _S74 + mat2x2<f32>(_S85[0] * _S84[0], _S85[1] * _S84[1]);
var vel_grad_det_3 : f32 = vel_grad_det_2 + _S82 * dot(_S83, _S78);
var _S87 : u32 = _S47 + u32(20);
var _S88 : vec3<f32> = shared_nodes_vel_mass_0[_S87];
var _S89 : NodeCdf_0 = shared_nodes_cdf_0[_S87];
var _S90 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(0), u32(2))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S87].affinities_0))
{
if((_S89.closest_id_0) != u32(4294967295))
{
var _S91 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S89.closest_id_0];
var _S92 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S89.closest_id_0].com_0, &(_S91), _S90 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S92 + project_velocity_0(_S41 - _S92, _S42.normal_0), _S88.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S88.z);
}
}
else
{
momentum_velocity_mass_1 = _S88;
}
var _S93 : f32 = _S45[i32(0)][u32(0)] * _S45[i32(1)][u32(2)];
var momentum_velocity_mass_4 : vec3<f32> = momentum_velocity_mass_3 + momentum_velocity_mass_1 * vec3<f32>(_S93);
var _S94 : f32 = _S93 * _S43;
var _S95 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S96 : mat2x2<f32> = outer_product_0(_S95, _S90);
var _S97 : mat2x2<f32> = mat2x2<f32>(_S94, _S94, _S94, _S94);
var _S98 : mat2x2<f32> = _S86 + mat2x2<f32>(_S97[0] * _S96[0], _S97[1] * _S96[1]);
var vel_grad_det_4 : f32 = vel_grad_det_3 + _S94 * dot(_S95, _S90);
var _S99 : vec3<f32> = shared_nodes_vel_mass_0[_S47];
var _S100 : NodeCdf_0 = shared_nodes_cdf_0[_S47];
var _S101 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(0), u32(0))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S47].affinities_0))
{
if((_S100.closest_id_0) != u32(4294967295))
{
var _S102 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S100.closest_id_0];
var _S103 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S100.closest_id_0].com_0, &(_S102), _S101 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S103 + project_velocity_0(_S41 - _S103, _S42.normal_0), _S99.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S99.z);
}
}
else
{
momentum_velocity_mass_1 = _S99;
}
var _S104 : f32 = _S45[i32(0)][u32(0)] * _S45[i32(1)][u32(0)];
var momentum_velocity_mass_5 : vec3<f32> = momentum_velocity_mass_4 + momentum_velocity_mass_1 * vec3<f32>(_S104);
var _S105 : f32 = _S104 * _S43;
var _S106 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S107 : mat2x2<f32> = outer_product_0(_S106, _S101);
var _S108 : mat2x2<f32> = mat2x2<f32>(_S105, _S105, _S105, _S105);
var _S109 : mat2x2<f32> = _S98 + mat2x2<f32>(_S108[0] * _S107[0], _S108[1] * _S107[1]);
var vel_grad_det_5 : f32 = vel_grad_det_4 + _S105 * dot(_S106, _S101);
var _S110 : u32 = _S47 + u32(10);
var _S111 : vec3<f32> = shared_nodes_vel_mass_0[_S110];
var _S112 : NodeCdf_0 = shared_nodes_cdf_0[_S110];
var _S113 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(0), u32(1))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S110].affinities_0))
{
if((_S112.closest_id_0) != u32(4294967295))
{
var _S114 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S112.closest_id_0];
var _S115 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S112.closest_id_0].com_0, &(_S114), _S113 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S115 + project_velocity_0(_S41 - _S115, _S42.normal_0), _S111.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S111.z);
}
}
else
{
momentum_velocity_mass_1 = _S111;
}
var _S116 : f32 = _S45[i32(0)][u32(0)] * _S45[i32(1)][u32(1)];
var momentum_velocity_mass_6 : vec3<f32> = momentum_velocity_mass_5 + momentum_velocity_mass_1 * vec3<f32>(_S116);
var _S117 : f32 = _S116 * _S43;
var _S118 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S119 : mat2x2<f32> = outer_product_0(_S118, _S113);
var _S120 : mat2x2<f32> = mat2x2<f32>(_S117, _S117, _S117, _S117);
var _S121 : mat2x2<f32> = _S109 + mat2x2<f32>(_S120[0] * _S119[0], _S120[1] * _S119[1]);
var vel_grad_det_6 : f32 = vel_grad_det_5 + _S117 * dot(_S118, _S113);
var _S122 : u32 = _S47 + u32(21);
var _S123 : vec3<f32> = shared_nodes_vel_mass_0[_S122];
var _S124 : NodeCdf_0 = shared_nodes_cdf_0[_S122];
var _S125 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(1), u32(2))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S122].affinities_0))
{
if((_S124.closest_id_0) != u32(4294967295))
{
var _S126 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S124.closest_id_0];
var _S127 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S124.closest_id_0].com_0, &(_S126), _S125 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S127 + project_velocity_0(_S41 - _S127, _S42.normal_0), _S123.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S123.z);
}
}
else
{
momentum_velocity_mass_1 = _S123;
}
var _S128 : f32 = _S45[i32(0)][u32(1)] * _S45[i32(1)][u32(2)];
var momentum_velocity_mass_7 : vec3<f32> = momentum_velocity_mass_6 + momentum_velocity_mass_1 * vec3<f32>(_S128);
var _S129 : f32 = _S128 * _S43;
var _S130 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S131 : mat2x2<f32> = outer_product_0(_S130, _S125);
var _S132 : mat2x2<f32> = mat2x2<f32>(_S129, _S129, _S129, _S129);
var _S133 : mat2x2<f32> = _S121 + mat2x2<f32>(_S132[0] * _S131[0], _S132[1] * _S131[1]);
var vel_grad_det_7 : f32 = vel_grad_det_6 + _S129 * dot(_S130, _S125);
var _S134 : u32 = _S47 + u32(1);
var _S135 : vec3<f32> = shared_nodes_vel_mass_0[_S134];
var _S136 : NodeCdf_0 = shared_nodes_cdf_0[_S134];
var _S137 : vec2<f32> = _S44 + vec2<f32>(vec2<u32>(u32(1), u32(0))) * _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S134].affinities_0))
{
if((_S136.closest_id_0) != u32(4294967295))
{
var _S138 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S136.closest_id_0];
var _S139 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S136.closest_id_0].com_0, &(_S138), _S137 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S139 + project_velocity_0(_S41 - _S139, _S42.normal_0), _S135.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S135.z);
}
}
else
{
momentum_velocity_mass_1 = _S135;
}
var _S140 : f32 = _S45[i32(0)][u32(1)] * _S45[i32(1)][u32(0)];
var momentum_velocity_mass_8 : vec3<f32> = momentum_velocity_mass_7 + momentum_velocity_mass_1 * vec3<f32>(_S140);
var _S141 : f32 = _S140 * _S43;
var _S142 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S143 : mat2x2<f32> = outer_product_0(_S142, _S137);
var _S144 : mat2x2<f32> = mat2x2<f32>(_S141, _S141, _S141, _S141);
var _S145 : mat2x2<f32> = _S133 + mat2x2<f32>(_S144[0] * _S143[0], _S144[1] * _S143[1]);
var vel_grad_det_8 : f32 = vel_grad_det_7 + _S141 * dot(_S142, _S137);
var _S146 : u32 = _S47 + u32(11);
var _S147 : vec3<f32> = shared_nodes_vel_mass_0[_S146];
var _S148 : NodeCdf_0 = shared_nodes_cdf_0[_S146];
var _S149 : vec2<f32> = _S44 + _S51;
if(!affinities_are_compatible_0(_S42.affinity_0, shared_nodes_cdf_0[_S146].affinities_0))
{
if((_S148.closest_id_0) != u32(4294967295))
{
var _S150 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S148.closest_id_0];
var _S151 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S148.closest_id_0].com_0, &(_S150), _S149 + _S40.pt_0);
momentum_velocity_mass_1 = vec3<f32>(_S151 + project_velocity_0(_S41 - _S151, _S42.normal_0), _S147.z);
}
else
{
momentum_velocity_mass_1 = vec3<f32>(_S41, _S147.z);
}
}
else
{
momentum_velocity_mass_1 = _S147;
}
var _S152 : f32 = _S45[i32(0)][u32(1)] * _S45[i32(1)][u32(1)];
var momentum_velocity_mass_9 : vec3<f32> = momentum_velocity_mass_8 + momentum_velocity_mass_1 * vec3<f32>(_S152);
var _S153 : f32 = _S152 * _S43;
var _S154 : vec2<f32> = momentum_velocity_mass_1.xy;
var _S155 : mat2x2<f32> = outer_product_0(_S154, _S149);
var _S156 : mat2x2<f32> = mat2x2<f32>(_S153, _S153, _S153, _S153);
var _S157 : mat2x2<f32> = _S145 + mat2x2<f32>(_S156[0] * _S155[0], _S156[1] * _S155[1]);
var vel_grad_det_9 : f32 = vel_grad_det_8 + _S153 * dot(_S154, _S149);
var i_5 : u32 = u32(0);
rigid_vel_1 = _S37;
for(;;)
{
if(i_5 < u32(16))
{
}
else
{
break;
}
if(affinity_bit_0(i_5, _S42.affinity_0))
{
var _S158 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[i_5];
var _S159 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[i_5].com_0, &(_S158), _S40.pt_0);
rigid_vel_1 = rigid_vel_1 + _S159;
}
i_5 = i_5 + u32(1);
}
velocity_gradient_0 = _S157;
vel_grad_det_1 = vel_grad_det_9;
momentum_velocity_mass_1 = momentum_velocity_mass_9;
}
else
{
rigid_vel_1 = _S37;
velocity_gradient_0 = _S39;
vel_grad_det_1 = 0.0f;
momentum_velocity_mass_1 = _S38;
}
entryPointParams_particles_dyn_0[_S34].cdf_1.rigid_vel_0 = rigid_vel_1;
var _S160 : array<vec2<f32>, i32(2)> = array<vec2<f32>, i32(2)>( vec2<f32>(velocity_gradient_0[i32(0)][i32(0)], velocity_gradient_0[i32(0)][i32(1)]), vec2<f32>(velocity_gradient_0[i32(1)][i32(0)], velocity_gradient_0[i32(1)][i32(1)]) );
entryPointParams_particles_dyn_0[_S34].affine_0.data_0 = _S160;
entryPointParams_particles_dyn_0[_S34].vel_grad_det_0 = vel_grad_det_1;
entryPointParams_particles_dyn_0[_S34].velocity_0 = momentum_velocity_mass_1.xy;
return;
}
@compute
@workgroup_size(8, 8, 1)
fn g2p(@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 _S161 : u32 = block_id_0.x;
var _S162 : ActiveBlockHeaderGeneric_std430_0 = entryPointParams_active_blocks_0[_S161];
global_shared_memory_transfers_0(tid_0, &(_S162.virtual_id_0));
workgroupBarrier();
var _S163 : u32 = entryPointParams_active_blocks_0[_S161].first_particle_0 + entryPointParams_active_blocks_0[_S161].num_particles_0;
var sorted_particle_id_0 : u32 = entryPointParams_active_blocks_0[_S161].first_particle_0 + tid_flat_0;
for(;;)
{
if(sorted_particle_id_0 < _S163)
{
}
else
{
break;
}
particle_g2p_0(entryPointParams_sorted_particle_ids_0[sorted_particle_id_0], entryPointParams_grid_0[i32(0)].cell_width_0, entryPointParams_params_0.dt_0);
sorted_particle_id_0 = sorted_particle_id_0 + u32(64);
}
return;
}