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_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) vel_grad_det_0 : f32,
@align(16) cdf_0 : Cdf_std430_0,
@align(16) init_volume_0 : f32,
@align(4) init_radius_0 : f32,
@align(8) mass_0 : f32,
@align(4) phase_0 : f32,
@align(16) 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 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_0 : ElasticitySoundSpeedTimestepBound_0, density0_0 : f32, def_grad_det_0 : f32, velocity_1 : vec3<f32>, cell_width_1 : f32) -> f32
{
return this_0.alpha_0 * cell_width_1 / max(length(velocity_1), sqrt((this_0.bulk_modulus_0 + this_0.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_1 : LinearElasticModel_0, particle_density0_0 : f32, particle_velocity_0 : vec3<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_1.cfl_coeff_0, bulk_modulus_from_lame_0(this_1.lambda_1, this_1.mu_2) * elastic_hardening_0, shear_modulus_from_lame_0(this_1.lambda_1, this_1.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_2 : NeoHookeanModel_0, particle_density0_1 : f32, particle_velocity_1 : vec3<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_2.cfl_coeff_1, bulk_modulus_from_lame_0(this_2.lambda_2, this_2.mu_3) * elastic_hardening_1, shear_modulus_from_lame_0(this_2.lambda_2, this_2.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));
}
fn Dynamics_init_density_0( _S2 : u32) -> f32
{
return entryPointParams_particles_dyn_0[_S2].mass_0 / entryPointParams_particles_dyn_0[_S2].init_volume_0;
}
struct DefaultParticleModel_0
{
dummy_0 : u32,
};
fn DefaultParticleModel_timestep_bound_0( _S3 : DefaultParticleModel_0, _S4 : u32, _S5 : f32, _S6 : mat3x3<f32>, _S7 : vec3<f32>, _S8 : f32) -> f32
{
var _S9 : u32 = _S4 * u32(52);
var _S10 : u32 = _S9 + u32(4);
var _S11 : u32 = entryPointParams_particles_model_0[(u32(i32(_S9)))/4];
var _S12 : f32 = determinant(_S6);
switch(i32(_S11))
{
case i32(0):
{
var _S13 : u32 = entryPointParams_particles_model_0[(_S10)/4];
var _S14 : f32 = bitcast<f32>(_S13);
var _S15 : u32 = entryPointParams_particles_model_0[(_S10 + u32(4))/4];
var _S16 : f32 = bitcast<f32>(_S15);
var _S17 : u32 = entryPointParams_particles_model_0[(_S10 + u32(8))/4];
var _S18 : LinearElasticModel_0 = LinearElasticModel_0( _S14, _S16, bitcast<f32>(_S17) );
return LinearElasticModel_timestep_bound_0(_S18, _S5, _S7, _S12, 1.0f, _S8);
}
case i32(1):
{
var _S19 : u32 = entryPointParams_particles_model_0[(_S10)/4];
var _S20 : f32 = bitcast<f32>(_S19);
var _S21 : u32 = entryPointParams_particles_model_0[(_S10 + u32(4))/4];
var _S22 : f32 = bitcast<f32>(_S21);
var _S23 : u32 = entryPointParams_particles_model_0[(_S10 + u32(8))/4];
var _S24 : NeoHookeanModel_0 = NeoHookeanModel_0( _S20, _S22, bitcast<f32>(_S23) );
return NeoHookeanModel_timestep_bound_0(_S24, _S5, _S7, _S12, 1.0f, _S8);
}
case i32(2):
{
var _S25 : u32 = entryPointParams_particles_model_0[(_S10)/4];
var _S26 : u32 = entryPointParams_particles_model_0[(_S10 + u32(4))/4];
var _S27 : u32 = entryPointParams_particles_model_0[(_S10 + u32(8))/4];
var _S28 : u32 = entryPointParams_particles_model_0[(_S10 + u32(12))/4];
var _S29 : u32 = entryPointParams_particles_model_0[(_S10 + u32(16))/4];
var _S30 : u32 = entryPointParams_particles_model_0[(_S10 + u32(20))/4];
var _S31 : u32 = entryPointParams_particles_model_0[(_S10 + u32(24))/4];
var _S32 : u32 = entryPointParams_particles_model_0[(_S10 + u32(28))/4];
var _S33 : u32 = entryPointParams_particles_model_0[(_S10 + u32(32))/4];
var _S34 : u32 = entryPointParams_particles_model_0[(_S10 + u32(36))/4];
var _S35 : f32 = bitcast<f32>(_S34);
var _S36 : u32 = entryPointParams_particles_model_0[(_S10 + u32(40))/4];
var _S37 : f32 = bitcast<f32>(_S36);
var _S38 : u32 = entryPointParams_particles_model_0[(_S10 + u32(44))/4];
var _S39 : LinearElasticModel_0 = LinearElasticModel_0( _S35, _S37, bitcast<f32>(_S38) );
return LinearElasticModel_timestep_bound_0(_S39, _S5, _S7, _S12, 1.0f, _S8);
}
case i32(3):
{
var _S40 : u32 = entryPointParams_particles_model_0[(_S10)/4];
var _S41 : u32 = entryPointParams_particles_model_0[(_S10 + u32(4))/4];
var _S42 : u32 = entryPointParams_particles_model_0[(_S10 + u32(8))/4];
var _S43 : u32 = entryPointParams_particles_model_0[(_S10 + u32(12))/4];
var _S44 : u32 = entryPointParams_particles_model_0[(_S10 + u32(16))/4];
var _S45 : u32 = entryPointParams_particles_model_0[(_S10 + u32(20))/4];
var _S46 : u32 = entryPointParams_particles_model_0[(_S10 + u32(24))/4];
var _S47 : u32 = entryPointParams_particles_model_0[(_S10 + u32(28))/4];
var _S48 : u32 = entryPointParams_particles_model_0[(_S10 + u32(32))/4];
var _S49 : u32 = entryPointParams_particles_model_0[(_S10 + u32(36))/4];
var _S50 : f32 = bitcast<f32>(_S49);
var _S51 : u32 = entryPointParams_particles_model_0[(_S10 + u32(40))/4];
var _S52 : f32 = bitcast<f32>(_S51);
var _S53 : u32 = entryPointParams_particles_model_0[(_S10 + u32(44))/4];
var _S54 : NeoHookeanModel_0 = NeoHookeanModel_0( _S50, _S52, bitcast<f32>(_S53) );
return NeoHookeanModel_timestep_bound_0(_S54, _S5, _S7, _S12, 1.0f, _S8);
}
default :
{
return 0.0f;
}
}
}
@compute
@workgroup_size(64, 1, 1)
fn estimate_timestep_bound(@builtin(global_invocation_id) invocation_id_0 : vec3<u32>)
{
var _S55 : u32 = invocation_id_0.x;
if(_S55 >= entryPointParams_particles_len_0)
{
return;
}
if((entryPointParams_particles_dyn_0[_S55].enabled_0) == u32(0))
{
return;
}
var _S56 : GridGeneric_std430_0 = entryPointParams_grid_0[i32(0)];
var _S57 : vec3<f32> = entryPointParams_particles_dyn_0[_S55].velocity_0;
var _S58 : mat3x3<f32> = mat3x3<f32>(entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(0)][i32(0)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(0)][i32(1)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(0)][i32(2)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(1)][i32(0)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(1)][i32(1)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(1)][i32(2)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(2)][i32(0)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(2)][i32(1)], entryPointParams_particles_dyn_0[_S55].affine_0.data_0[i32(2)][i32(2)]);
var _S59 : f32 = entryPointParams_particles_dyn_0[_S55].mass_0;
var _S60 : DefaultParticleModel_0 = DefaultParticleModel_0( );
var _S61 : f32 = DefaultParticleModel_timestep_bound_0(_S60, _S55, Dynamics_init_density_0(_S55), mat3x3<f32>(entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(0)][i32(0)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(0)][i32(1)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(0)][i32(2)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(1)][i32(0)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(1)][i32(1)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(1)][i32(2)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(2)][i32(0)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(2)][i32(1)], entryPointParams_particles_dyn_0[_S55].def_grad_0.data_0[i32(2)][i32(2)]), entryPointParams_particles_dyn_0[_S55].velocity_0, _S56.cell_width_0);
var _S62 : f32 = _S58[i32(0)][i32(0)];
var _S63 : f32 = _S58[i32(0)][i32(1)];
var _S64 : f32 = _S58[i32(0)][i32(2)];
var _S65 : f32 = _S58[i32(1)][i32(0)];
var _S66 : f32 = _S58[i32(1)][i32(1)];
var _S67 : f32 = _S58[i32(1)][i32(2)];
var _S68 : f32 = _S58[i32(2)][i32(0)];
var _S69 : f32 = _S58[i32(2)][i32(1)];
var _S70 : f32 = _S58[i32(2)][i32(2)];
var _S71 : u32 = atomicMin(&(entryPointParams_result_1[i32(0)].computed_max_dt_as_uint_0), GpuTimestepBounds_secs_to_int_0(min(_S61, _S56.cell_width_0 / (length(_S57) + _S56.cell_width_0 * _S56.cell_width_0 / 4.0f * sqrt(_S62 * _S62 + _S63 * _S63 + _S64 * _S64 + _S65 * _S65 + _S66 * _S66 + _S67 * _S67 + _S68 * _S68 + _S69 * _S69 + _S70 * _S70) / _S59 * 6.0f * sqrt(3.0f) / _S56.cell_width_0))));
return;
}