struct GpuTimestepBounds_std430_0
{
@align(4) computed_max_dt_as_uint_0 : atomic<u32>,
};
@binding(0) @group(0) var<storage, read_write> entryPointParams_result_0 : array<GpuTimestepBounds_std430_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>;
@binding(2) @group(0) var<storage, read> entryPointParams_particles_model_0 : array<u32>;
struct _MatrixStorage_float2x2std430_0
{
@align(8) data_0 : array<vec2<f32>, i32(2)>,
};
struct Cdf_std430_0
{
@align(8) normal_0 : vec2<f32>,
@align(8) rigid_vel_0 : vec2<f32>,
@align(8) signed_distance_0 : f32,
@align(4) affinity_0 : u32,
};
struct Dynamics_std430_0
{
@align(8) velocity_0 : vec2<f32>,
@align(8) def_grad_0 : _MatrixStorage_float2x2std430_0,
@align(8) affine_0 : _MatrixStorage_float2x2std430_0,
@align(8) vel_grad_det_0 : f32,
@align(8) cdf_0 : Cdf_std430_0,
@align(8) init_volume_0 : f32,
@align(4) init_radius_0 : f32,
@align(8) mass_0 : f32,
@align(4) phase_0 : f32,
@align(8) enabled_0 : u32,
};
@binding(3) @group(0) var<storage, read> entryPointParams_particles_dyn_0 : array<Dynamics_std430_0>;
@binding(4) @group(0) var<uniform> entryPointParams_particles_len_0 : u32;
@binding(5) @group(0) var<storage, read_write> entryPointParams_result_1 : array<GpuTimestepBounds_std430_0>;
@compute
@workgroup_size(1, 1, 1)
fn reset_timestep_bound()
{
atomicStore(&(entryPointParams_result_0[i32(0)].computed_max_dt_as_uint_0), u32(4294967295));
return;
}
fn Dynamics_init_density_0( this_0 : ptr<function, Dynamics_std430_0>) -> f32
{
return (*this_0).mass_0 / (*this_0).init_volume_0;
}
fn bulk_modulus_from_lame_0( lambda_0 : f32, mu_0 : f32) -> f32
{
return lambda_0 + 2.0f * mu_0 / 3.0f;
}
fn shear_modulus_from_lame_0( _lambda_0 : f32, mu_1 : f32) -> f32
{
return mu_1;
}
fn hook_to_bulk_modulus_0( young_modulus_0 : f32, poisson_ratio_0 : f32) -> f32
{
return young_modulus_0 / (3.0f * (1.0f - 2.0f * poisson_ratio_0));
}
fn hook_to_shear_modulus_0( young_modulus_1 : f32, poisson_ratio_1 : f32) -> f32
{
return young_modulus_1 / (2.0f * (1.0f + poisson_ratio_1));
}
struct ElasticitySoundSpeedTimestepBound_0
{
alpha_0 : f32,
bulk_modulus_0 : f32,
shear_modulus_0 : f32,
};
fn ElasticitySoundSpeedTimestepBound_x24init_0( alpha_1 : f32, young_modulus_2 : f32, poisson_ratio_2 : f32) -> ElasticitySoundSpeedTimestepBound_0
{
var _S1 : ElasticitySoundSpeedTimestepBound_0;
_S1.alpha_0 = alpha_1;
_S1.bulk_modulus_0 = hook_to_bulk_modulus_0(young_modulus_2, poisson_ratio_2);
_S1.shear_modulus_0 = hook_to_shear_modulus_0(young_modulus_2, poisson_ratio_2);
return _S1;
}
fn ElasticitySoundSpeedTimestepBound_timestep_bound_0( this_1 : ElasticitySoundSpeedTimestepBound_0, density0_0 : f32, def_grad_det_0 : f32, velocity_1 : vec2<f32>, cell_width_1 : f32) -> f32
{
return this_1.alpha_0 * cell_width_1 / max(length(velocity_1), sqrt((this_1.bulk_modulus_0 + this_1.shear_modulus_0 * 4.0f / 3.0f) / (density0_0 / max(def_grad_det_0, 9.99999997475242708e-07f))));
}
struct LinearElasticModel_0
{
lambda_1 : f32,
mu_2 : f32,
cfl_coeff_0 : f32,
};
fn LinearElasticModel_timestep_bound_0( this_2 : LinearElasticModel_0, particle_density0_0 : f32, particle_velocity_0 : vec2<f32>, particle_def_grad_det_0 : f32, elastic_hardening_0 : f32, cell_width_2 : f32) -> f32
{
return ElasticitySoundSpeedTimestepBound_timestep_bound_0(ElasticitySoundSpeedTimestepBound_x24init_0(this_2.cfl_coeff_0, bulk_modulus_from_lame_0(this_2.lambda_1, this_2.mu_2) * elastic_hardening_0, shear_modulus_from_lame_0(this_2.lambda_1, this_2.mu_2) * elastic_hardening_0), particle_density0_0, particle_def_grad_det_0, particle_velocity_0, cell_width_2);
}
struct NeoHookeanModel_0
{
lambda_2 : f32,
mu_3 : f32,
cfl_coeff_1 : f32,
};
fn NeoHookeanModel_timestep_bound_0( this_3 : NeoHookeanModel_0, particle_density0_1 : f32, particle_velocity_1 : vec2<f32>, particle_def_grad_det_1 : f32, elastic_hardening_1 : f32, cell_width_3 : f32) -> f32
{
return ElasticitySoundSpeedTimestepBound_timestep_bound_0(ElasticitySoundSpeedTimestepBound_x24init_0(this_3.cfl_coeff_1, bulk_modulus_from_lame_0(this_3.lambda_2, this_3.mu_3) * elastic_hardening_1, shear_modulus_from_lame_0(this_3.lambda_2, this_3.mu_3) * elastic_hardening_1), particle_density0_1, particle_def_grad_det_1, particle_velocity_1, cell_width_3);
}
fn GpuTimestepBounds_secs_to_int_0( secs_0 : f32) -> u32
{
return u32(floor(secs_0 * 9.99999995904e+11f));
}
struct DefaultParticleModel_0
{
dummy_0 : u32,
};
fn DefaultParticleModel_timestep_bound_0( _S2 : DefaultParticleModel_0, _S3 : u32, _S4 : f32, _S5 : mat2x2<f32>, _S6 : vec2<f32>, _S7 : f32) -> f32
{
var _S8 : u32 = _S3 * u32(52);
var _S9 : u32 = _S8 + u32(4);
var _S10 : u32 = entryPointParams_particles_model_0[(u32(i32(_S8)))/4];
var _S11 : f32 = determinant(_S5);
switch(i32(_S10))
{
case i32(0):
{
var _S12 : u32 = entryPointParams_particles_model_0[(_S9)/4];
var _S13 : f32 = bitcast<f32>(_S12);
var _S14 : u32 = entryPointParams_particles_model_0[(_S9 + u32(4))/4];
var _S15 : f32 = bitcast<f32>(_S14);
var _S16 : u32 = entryPointParams_particles_model_0[(_S9 + u32(8))/4];
var _S17 : LinearElasticModel_0 = LinearElasticModel_0( _S13, _S15, bitcast<f32>(_S16) );
return LinearElasticModel_timestep_bound_0(_S17, _S4, _S6, _S11, 1.0f, _S7);
}
case i32(1):
{
var _S18 : u32 = entryPointParams_particles_model_0[(_S9)/4];
var _S19 : f32 = bitcast<f32>(_S18);
var _S20 : u32 = entryPointParams_particles_model_0[(_S9 + u32(4))/4];
var _S21 : f32 = bitcast<f32>(_S20);
var _S22 : u32 = entryPointParams_particles_model_0[(_S9 + u32(8))/4];
var _S23 : NeoHookeanModel_0 = NeoHookeanModel_0( _S19, _S21, bitcast<f32>(_S22) );
return NeoHookeanModel_timestep_bound_0(_S23, _S4, _S6, _S11, 1.0f, _S7);
}
case i32(2):
{
var _S24 : u32 = entryPointParams_particles_model_0[(_S9)/4];
var _S25 : u32 = entryPointParams_particles_model_0[(_S9 + u32(4))/4];
var _S26 : u32 = entryPointParams_particles_model_0[(_S9 + u32(8))/4];
var _S27 : u32 = entryPointParams_particles_model_0[(_S9 + u32(12))/4];
var _S28 : u32 = entryPointParams_particles_model_0[(_S9 + u32(16))/4];
var _S29 : u32 = entryPointParams_particles_model_0[(_S9 + u32(20))/4];
var _S30 : u32 = entryPointParams_particles_model_0[(_S9 + u32(24))/4];
var _S31 : u32 = entryPointParams_particles_model_0[(_S9 + u32(28))/4];
var _S32 : u32 = entryPointParams_particles_model_0[(_S9 + u32(32))/4];
var _S33 : u32 = entryPointParams_particles_model_0[(_S9 + u32(36))/4];
var _S34 : f32 = bitcast<f32>(_S33);
var _S35 : u32 = entryPointParams_particles_model_0[(_S9 + u32(40))/4];
var _S36 : f32 = bitcast<f32>(_S35);
var _S37 : u32 = entryPointParams_particles_model_0[(_S9 + u32(44))/4];
var _S38 : LinearElasticModel_0 = LinearElasticModel_0( _S34, _S36, bitcast<f32>(_S37) );
return LinearElasticModel_timestep_bound_0(_S38, _S4, _S6, _S11, 1.0f, _S7);
}
case i32(3):
{
var _S39 : u32 = entryPointParams_particles_model_0[(_S9)/4];
var _S40 : u32 = entryPointParams_particles_model_0[(_S9 + u32(4))/4];
var _S41 : u32 = entryPointParams_particles_model_0[(_S9 + u32(8))/4];
var _S42 : u32 = entryPointParams_particles_model_0[(_S9 + u32(12))/4];
var _S43 : u32 = entryPointParams_particles_model_0[(_S9 + u32(16))/4];
var _S44 : u32 = entryPointParams_particles_model_0[(_S9 + u32(20))/4];
var _S45 : u32 = entryPointParams_particles_model_0[(_S9 + u32(24))/4];
var _S46 : u32 = entryPointParams_particles_model_0[(_S9 + u32(28))/4];
var _S47 : u32 = entryPointParams_particles_model_0[(_S9 + u32(32))/4];
var _S48 : u32 = entryPointParams_particles_model_0[(_S9 + u32(36))/4];
var _S49 : f32 = bitcast<f32>(_S48);
var _S50 : u32 = entryPointParams_particles_model_0[(_S9 + u32(40))/4];
var _S51 : f32 = bitcast<f32>(_S50);
var _S52 : u32 = entryPointParams_particles_model_0[(_S9 + u32(44))/4];
var _S53 : NeoHookeanModel_0 = NeoHookeanModel_0( _S49, _S51, bitcast<f32>(_S52) );
return NeoHookeanModel_timestep_bound_0(_S53, _S4, _S6, _S11, 1.0f, _S7);
}
default :
{
return 0.0f;
}
}
}
@compute
@workgroup_size(64, 1, 1)
fn estimate_timestep_bound(@builtin(global_invocation_id) invocation_id_0 : vec3<u32>)
{
var _S54 : u32 = invocation_id_0.x;
if(_S54 >= entryPointParams_particles_len_0)
{
return;
}
if((entryPointParams_particles_dyn_0[_S54].enabled_0) == u32(0))
{
return;
}
var _S55 : GridGeneric_std430_0 = entryPointParams_grid_0[i32(0)];
var _S56 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S54];
var _S57 : f32 = Dynamics_init_density_0(&(_S56));
var _S58 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S54];
var _S59 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S54];
var _S60 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S54];
var _S61 : mat2x2<f32> = mat2x2<f32>(_S60.affine_0.data_0[i32(0)][i32(0)], _S60.affine_0.data_0[i32(0)][i32(1)], _S60.affine_0.data_0[i32(1)][i32(0)], _S60.affine_0.data_0[i32(1)][i32(1)]);
var _S62 : Dynamics_std430_0 = entryPointParams_particles_dyn_0[_S54];
var _S63 : DefaultParticleModel_0 = DefaultParticleModel_0( );
var _S64 : f32 = DefaultParticleModel_timestep_bound_0(_S63, _S54, _S57, mat2x2<f32>(_S58.def_grad_0.data_0[i32(0)][i32(0)], _S58.def_grad_0.data_0[i32(0)][i32(1)], _S58.def_grad_0.data_0[i32(1)][i32(0)], _S58.def_grad_0.data_0[i32(1)][i32(1)]), _S59.velocity_0, _S55.cell_width_0);
var _S65 : f32 = _S61[i32(0)][i32(0)];
var _S66 : f32 = _S61[i32(0)][i32(1)];
var _S67 : f32 = _S61[i32(1)][i32(0)];
var _S68 : f32 = _S61[i32(1)][i32(1)];
var _S69 : u32 = atomicMin(&(entryPointParams_result_1[i32(0)].computed_max_dt_as_uint_0), GpuTimestepBounds_secs_to_int_0(min(_S64, _S55.cell_width_0 / (length(_S59.velocity_0) + _S55.cell_width_0 * _S55.cell_width_0 / 4.0f * sqrt(_S65 * _S65 + _S66 * _S66 + _S67 * _S67 + _S68 * _S68) / _S62.mass_0 * 6.0f * sqrt(2.0f) / _S55.cell_width_0))));
return;
}