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) cdf_1 : Cdf_std430_0,
@align(8) init_volume_0 : f32,
@align(4) init_radius_0 : f32,
@align(8) mass_0 : f32,
};
@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
{
@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;
}
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
{
@align(4) distance_0 : f32,
@align(4) affinities_0 : u32,
@align(4) 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);
}
struct Position_0
{
@align(8) pt_0 : vec2<f32>,
};
fn associated_grid_pos_0( part_pos_0 : Position_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 : Position_0, cell_width_3 : f32) -> vec2<f32>
{
return associated_grid_pos_0(part_pos_1, cell_width_3) - part_pos_1.pt_0;
}
fn QuadraticKernel_eval_all_0( x_1 : f32) -> vec3<f32>
{
var _S7 : f32 = 1.5f - x_1;
var _S8 : f32 = x_1 - 1.0f;
var _S9 : f32 = x_1 - 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 associated_cell_index_in_block_off_by_one_0( part_pos_2 : Position_0, cell_width_4 : f32) -> vec2<u32>
{
var _S10 : vec2<f32> = round(part_pos_2.pt_0 / vec2<f32>(cell_width_4)) - vec2<f32>(1.0f);
var _S11 : vec2<f32> = vec2<f32>(8.0f);
return vec2<u32>(_S10 - floor(_S10 / _S11) * _S11);
}
fn affinities_are_compatible_0( affinity1_0 : u32, affinity2_0 : u32) -> bool
{
var _S12 : u32 = (((affinity1_0 & (affinity2_0))) & (u32(65535)));
return ((((affinity1_0 >> (u32(16)))) & (_S12))) == ((((affinity2_0 >> (u32(16)))) & (_S12)));
}
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 _S13 : vec2<f32> = point_0 - center_of_mass_0;
return vels_0.linear_0 + vec2<f32>(vels_0.angular_0) * vec2<f32>(- _S13.y, _S13.x);
}
fn project_velocity_0( vel_0 : vec2<f32>, n_0 : vec2<f32>) -> vec2<f32>
{
var _S14 : f32 = dot(vel_0, n_0);
if(_S14 < 0.0f)
{
var _S15 : vec2<f32> = vel_0 - n_0 * vec2<f32>(_S14);
var _S16 : f32 = length(_S15);
return select(vec2<f32>(0.0f), _S15 / vec2<f32>(_S16), _S16 > 9.99999993922529029e-09f) * vec2<f32>(max(0.0f, _S16 + 20.0f * _S14));
}
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( _S17 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S18 : u32 = pack_key_0(_S17);
var slot_0 : u32 = ((hash_0(_S18)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S19 : u32 = entryPointParams_hmap_entries_0[slot_0].state_0;
if((entryPointParams_hmap_entries_0[slot_0].state_0) == _S18)
{
var _S20 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S20;
}
else
{
if(_S19 == 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( _S21 : vec3<u32>, _S22 : BlockVirtualId_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 _S23 : bool;
if(i_4 == u32(1))
{
_S23 = (_S21.x) > u32(1);
}
else
{
_S23 = false;
}
var _S24 : bool;
if(_S23)
{
_S24 = true;
}
else
{
if(j_0 == u32(1))
{
_S24 = (_S21.y) > u32(1);
}
else
{
_S24 = false;
}
}
if(_S24)
{
j_0 = j_0 + u32(1);
continue;
}
var _S25 : vec2<u32> = vec2<u32>(i_4, j_0);
var _S26 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S22.id_0 + vec2<i32>(_S25)));
var _S27 : vec2<u32> = _S21.xy;
var _S28 : vec2<u32> = _S25 * vec2<u32>(u32(8)) + _S27;
var _S29 : u32 = flatten_shared_index_0(_S28.x, _S28.y);
if((_S26.id_1) != u32(4294967295))
{
var _S30 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S26), _S27);
shared_nodes_vel_mass_0[_S29] = entryPointParams_nodes_0[_S30.id_3].momentum_velocity_mass_0;
var _S31 : u32 = entryPointParams_nodes_0[_S30.id_3].cdf_0.affinities_0;
var _S32 : u32 = entryPointParams_nodes_0[_S30.id_3].cdf_0.closest_id_0;
shared_nodes_cdf_0[_S29].distance_0 = entryPointParams_nodes_0[_S30.id_3].cdf_0.distance_0;
shared_nodes_cdf_0[_S29].affinities_0 = _S31;
shared_nodes_cdf_0[_S29].closest_id_0 = _S32;
}
else
{
shared_nodes_vel_mass_0[_S29] = vec3<f32>(0.0f);
shared_nodes_cdf_0[_S29] = 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( _S33 : u32, _S34 : f32, _S35 : f32)
{
var _S36 : vec2<f32> = vec2<f32>(0.0f);
var _S37 : Position_0 = Position_0( entryPointParams_particles_pos_0[_S33].pt_0 );
var _S38 : vec2<f32> = entryPointParams_particles_dyn_0[_S33].velocity_0;
var _S39 : vec2<f32> = entryPointParams_particles_dyn_0[_S33].cdf_1.normal_0;
var _S40 : u32 = entryPointParams_particles_dyn_0[_S33].cdf_1.affinity_0;
var _S41 : f32 = QuadraticKernel_inv_d_0(_S34);
var _S42 : vec2<f32> = dir_to_associated_grid_node_0(_S37, _S34);
var _S43 : mat2x3<f32> = QuadraticKernel_precompute_weights_0(_S42, _S34);
var _S44 : vec2<u32> = associated_cell_index_in_block_off_by_one_0(_S37, _S34);
var _S45 : u32 = flatten_shared_index_0(_S44.x, _S44.y);
var _S46 : u32 = _S45 + u32(22);
var _S47 : vec3<f32> = shared_nodes_vel_mass_0[_S46];
var _S48 : NodeCdf_0 = shared_nodes_cdf_0[_S46];
var _S49 : vec2<f32> = vec2<f32>(_S34);
var _S50 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(2), u32(2))) * _S49;
var cpic_cell_data_0 : vec3<f32>;
if(!affinities_are_compatible_0(entryPointParams_particles_dyn_0[_S33].cdf_1.affinity_0, shared_nodes_cdf_0[_S46].affinities_0))
{
if((_S48.closest_id_0) != u32(4294967295))
{
var _S51 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S48.closest_id_0].linear_0, entryPointParams_body_vels_0[_S48.closest_id_0].angular_0 );
var _S52 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S48.closest_id_0].com_0, _S51, _S50 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S52 + project_velocity_0(_S38 - _S52, _S39), _S47.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S47.z);
}
}
else
{
cpic_cell_data_0 = _S47;
}
var _S53 : f32 = _S43[i32(0)][u32(2)] * _S43[i32(1)][u32(2)];
var _S54 : vec3<f32> = cpic_cell_data_0 * vec3<f32>(_S53);
var _S55 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S50);
var _S56 : mat2x2<f32> = mat2x2<f32>(_S53 * _S41, _S53 * _S41, _S53 * _S41, _S53 * _S41);
var _S57 : mat2x2<f32> = mat2x2<f32>(_S56[0] * _S55[0], _S56[1] * _S55[1]);
var _S58 : u32 = _S45 + u32(2);
var _S59 : vec3<f32> = shared_nodes_vel_mass_0[_S58];
var _S60 : NodeCdf_0 = shared_nodes_cdf_0[_S58];
var _S61 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(2), u32(0))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S58].affinities_0))
{
if((_S60.closest_id_0) != u32(4294967295))
{
var _S62 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S60.closest_id_0].linear_0, entryPointParams_body_vels_0[_S60.closest_id_0].angular_0 );
var _S63 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S60.closest_id_0].com_0, _S62, _S61 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S63 + project_velocity_0(_S38 - _S63, _S39), _S59.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S59.z);
}
}
else
{
cpic_cell_data_0 = _S59;
}
var _S64 : f32 = _S43[i32(0)][u32(2)] * _S43[i32(1)][u32(0)];
var momentum_velocity_mass_1 : vec3<f32> = _S54 + cpic_cell_data_0 * vec3<f32>(_S64);
var _S65 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S61);
var _S66 : mat2x2<f32> = mat2x2<f32>(_S64 * _S41, _S64 * _S41, _S64 * _S41, _S64 * _S41);
var velocity_gradient_0 : mat2x2<f32> = _S57 + mat2x2<f32>(_S66[0] * _S65[0], _S66[1] * _S65[1]);
var _S67 : u32 = _S45 + u32(12);
var _S68 : vec3<f32> = shared_nodes_vel_mass_0[_S67];
var _S69 : NodeCdf_0 = shared_nodes_cdf_0[_S67];
var _S70 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(2), u32(1))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S67].affinities_0))
{
if((_S69.closest_id_0) != u32(4294967295))
{
var _S71 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S69.closest_id_0].linear_0, entryPointParams_body_vels_0[_S69.closest_id_0].angular_0 );
var _S72 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S69.closest_id_0].com_0, _S71, _S70 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S72 + project_velocity_0(_S38 - _S72, _S39), _S68.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S68.z);
}
}
else
{
cpic_cell_data_0 = _S68;
}
var _S73 : f32 = _S43[i32(0)][u32(2)] * _S43[i32(1)][u32(1)];
var momentum_velocity_mass_2 : vec3<f32> = momentum_velocity_mass_1 + cpic_cell_data_0 * vec3<f32>(_S73);
var _S74 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S70);
var _S75 : mat2x2<f32> = mat2x2<f32>(_S73 * _S41, _S73 * _S41, _S73 * _S41, _S73 * _S41);
var velocity_gradient_1 : mat2x2<f32> = velocity_gradient_0 + mat2x2<f32>(_S75[0] * _S74[0], _S75[1] * _S74[1]);
var _S76 : u32 = _S45 + u32(20);
var _S77 : vec3<f32> = shared_nodes_vel_mass_0[_S76];
var _S78 : NodeCdf_0 = shared_nodes_cdf_0[_S76];
var _S79 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(0), u32(2))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S76].affinities_0))
{
if((_S78.closest_id_0) != u32(4294967295))
{
var _S80 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S78.closest_id_0].linear_0, entryPointParams_body_vels_0[_S78.closest_id_0].angular_0 );
var _S81 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S78.closest_id_0].com_0, _S80, _S79 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S81 + project_velocity_0(_S38 - _S81, _S39), _S77.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S77.z);
}
}
else
{
cpic_cell_data_0 = _S77;
}
var _S82 : f32 = _S43[i32(0)][u32(0)] * _S43[i32(1)][u32(2)];
var momentum_velocity_mass_3 : vec3<f32> = momentum_velocity_mass_2 + cpic_cell_data_0 * vec3<f32>(_S82);
var _S83 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S79);
var _S84 : mat2x2<f32> = mat2x2<f32>(_S82 * _S41, _S82 * _S41, _S82 * _S41, _S82 * _S41);
var velocity_gradient_2 : mat2x2<f32> = velocity_gradient_1 + mat2x2<f32>(_S84[0] * _S83[0], _S84[1] * _S83[1]);
var _S85 : vec3<f32> = shared_nodes_vel_mass_0[_S45];
var _S86 : NodeCdf_0 = shared_nodes_cdf_0[_S45];
var _S87 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(0), u32(0))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S45].affinities_0))
{
if((_S86.closest_id_0) != u32(4294967295))
{
var _S88 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S86.closest_id_0].linear_0, entryPointParams_body_vels_0[_S86.closest_id_0].angular_0 );
var _S89 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S86.closest_id_0].com_0, _S88, _S87 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S89 + project_velocity_0(_S38 - _S89, _S39), _S85.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S85.z);
}
}
else
{
cpic_cell_data_0 = _S85;
}
var _S90 : f32 = _S43[i32(0)][u32(0)] * _S43[i32(1)][u32(0)];
var momentum_velocity_mass_4 : vec3<f32> = momentum_velocity_mass_3 + cpic_cell_data_0 * vec3<f32>(_S90);
var _S91 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S87);
var _S92 : mat2x2<f32> = mat2x2<f32>(_S90 * _S41, _S90 * _S41, _S90 * _S41, _S90 * _S41);
var velocity_gradient_3 : mat2x2<f32> = velocity_gradient_2 + mat2x2<f32>(_S92[0] * _S91[0], _S92[1] * _S91[1]);
var _S93 : u32 = _S45 + u32(10);
var _S94 : vec3<f32> = shared_nodes_vel_mass_0[_S93];
var _S95 : NodeCdf_0 = shared_nodes_cdf_0[_S93];
var _S96 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(0), u32(1))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S93].affinities_0))
{
if((_S95.closest_id_0) != u32(4294967295))
{
var _S97 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S95.closest_id_0].linear_0, entryPointParams_body_vels_0[_S95.closest_id_0].angular_0 );
var _S98 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S95.closest_id_0].com_0, _S97, _S96 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S98 + project_velocity_0(_S38 - _S98, _S39), _S94.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S94.z);
}
}
else
{
cpic_cell_data_0 = _S94;
}
var _S99 : f32 = _S43[i32(0)][u32(0)] * _S43[i32(1)][u32(1)];
var momentum_velocity_mass_5 : vec3<f32> = momentum_velocity_mass_4 + cpic_cell_data_0 * vec3<f32>(_S99);
var _S100 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S96);
var _S101 : mat2x2<f32> = mat2x2<f32>(_S99 * _S41, _S99 * _S41, _S99 * _S41, _S99 * _S41);
var velocity_gradient_4 : mat2x2<f32> = velocity_gradient_3 + mat2x2<f32>(_S101[0] * _S100[0], _S101[1] * _S100[1]);
var _S102 : u32 = _S45 + u32(21);
var _S103 : vec3<f32> = shared_nodes_vel_mass_0[_S102];
var _S104 : NodeCdf_0 = shared_nodes_cdf_0[_S102];
var _S105 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(1), u32(2))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S102].affinities_0))
{
if((_S104.closest_id_0) != u32(4294967295))
{
var _S106 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S104.closest_id_0].linear_0, entryPointParams_body_vels_0[_S104.closest_id_0].angular_0 );
var _S107 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S104.closest_id_0].com_0, _S106, _S105 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S107 + project_velocity_0(_S38 - _S107, _S39), _S103.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S103.z);
}
}
else
{
cpic_cell_data_0 = _S103;
}
var _S108 : f32 = _S43[i32(0)][u32(1)] * _S43[i32(1)][u32(2)];
var momentum_velocity_mass_6 : vec3<f32> = momentum_velocity_mass_5 + cpic_cell_data_0 * vec3<f32>(_S108);
var _S109 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S105);
var _S110 : mat2x2<f32> = mat2x2<f32>(_S108 * _S41, _S108 * _S41, _S108 * _S41, _S108 * _S41);
var velocity_gradient_5 : mat2x2<f32> = velocity_gradient_4 + mat2x2<f32>(_S110[0] * _S109[0], _S110[1] * _S109[1]);
var _S111 : u32 = _S45 + u32(1);
var _S112 : vec3<f32> = shared_nodes_vel_mass_0[_S111];
var _S113 : NodeCdf_0 = shared_nodes_cdf_0[_S111];
var _S114 : vec2<f32> = _S42 + vec2<f32>(vec2<u32>(u32(1), u32(0))) * _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S111].affinities_0))
{
if((_S113.closest_id_0) != u32(4294967295))
{
var _S115 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S113.closest_id_0].linear_0, entryPointParams_body_vels_0[_S113.closest_id_0].angular_0 );
var _S116 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S113.closest_id_0].com_0, _S115, _S114 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S116 + project_velocity_0(_S38 - _S116, _S39), _S112.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S112.z);
}
}
else
{
cpic_cell_data_0 = _S112;
}
var _S117 : f32 = _S43[i32(0)][u32(1)] * _S43[i32(1)][u32(0)];
var momentum_velocity_mass_7 : vec3<f32> = momentum_velocity_mass_6 + cpic_cell_data_0 * vec3<f32>(_S117);
var _S118 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S114);
var _S119 : mat2x2<f32> = mat2x2<f32>(_S117 * _S41, _S117 * _S41, _S117 * _S41, _S117 * _S41);
var velocity_gradient_6 : mat2x2<f32> = velocity_gradient_5 + mat2x2<f32>(_S119[0] * _S118[0], _S119[1] * _S118[1]);
var _S120 : u32 = _S45 + u32(11);
var _S121 : vec3<f32> = shared_nodes_vel_mass_0[_S120];
var _S122 : NodeCdf_0 = shared_nodes_cdf_0[_S120];
var _S123 : vec2<f32> = _S42 + _S49;
if(!affinities_are_compatible_0(_S40, shared_nodes_cdf_0[_S120].affinities_0))
{
if((_S122.closest_id_0) != u32(4294967295))
{
var _S124 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[_S122.closest_id_0].linear_0, entryPointParams_body_vels_0[_S122.closest_id_0].angular_0 );
var _S125 : vec2<f32> = body_velocity_at_point_0(entryPointParams_body_mprops_0[_S122.closest_id_0].com_0, _S124, _S123 + entryPointParams_particles_pos_0[_S33].pt_0);
cpic_cell_data_0 = vec3<f32>(_S125 + project_velocity_0(_S38 - _S125, _S39), _S121.z);
}
else
{
cpic_cell_data_0 = vec3<f32>(_S38, _S121.z);
}
}
else
{
cpic_cell_data_0 = _S121;
}
var _S126 : f32 = _S43[i32(0)][u32(1)] * _S43[i32(1)][u32(1)];
var momentum_velocity_mass_8 : vec3<f32> = momentum_velocity_mass_7 + cpic_cell_data_0 * vec3<f32>(_S126);
var _S127 : mat2x2<f32> = outer_product_0(cpic_cell_data_0.xy, _S123);
var _S128 : mat2x2<f32> = mat2x2<f32>(_S126 * _S41, _S126 * _S41, _S126 * _S41, _S126 * _S41);
var velocity_gradient_7 : mat2x2<f32> = velocity_gradient_6 + mat2x2<f32>(_S128[0] * _S127[0], _S128[1] * _S127[1]);
var i_5 : u32 = u32(0);
var rigid_vel_1 : vec2<f32> = _S36;
for(;;)
{
if(i_5 < u32(16))
{
}
else
{
break;
}
if(affinity_bit_0(i_5, _S40))
{
var _S129 : BodyVelocity_0 = BodyVelocity_0( entryPointParams_body_vels_0[i_5].linear_0, entryPointParams_body_vels_0[i_5].angular_0 );
rigid_vel_1 = rigid_vel_1 + body_velocity_at_point_0(entryPointParams_body_mprops_0[i_5].com_0, _S129, entryPointParams_particles_pos_0[_S33].pt_0);
}
i_5 = i_5 + u32(1);
}
entryPointParams_particles_dyn_0[_S33].cdf_1.rigid_vel_0 = rigid_vel_1;
var _S130 : array<vec2<f32>, i32(2)> = array<vec2<f32>, i32(2)>( vec2<f32>(velocity_gradient_7[i32(0)][i32(0)], velocity_gradient_7[i32(0)][i32(1)]), vec2<f32>(velocity_gradient_7[i32(1)][i32(0)], velocity_gradient_7[i32(1)][i32(1)]) );
entryPointParams_particles_dyn_0[_S33].affine_0.data_0 = _S130;
entryPointParams_particles_dyn_0[_S33].velocity_0 = momentum_velocity_mass_8.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 _S131 : u32 = block_id_0.x;
var _S132 : BlockVirtualId_0 = BlockVirtualId_0( entryPointParams_active_blocks_0[_S131].virtual_id_0.id_0 );
global_shared_memory_transfers_0(tid_0, _S132);
workgroupBarrier();
var _S133 : u32 = entryPointParams_active_blocks_0[_S131].first_particle_0 + entryPointParams_active_blocks_0[_S131].num_particles_0;
var sorted_particle_id_0 : u32 = entryPointParams_active_blocks_0[_S131].first_particle_0 + tid_flat_0;
for(;;)
{
if(sorted_particle_id_0 < _S133)
{
}
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;
}