slosh3d 0.6.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(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) 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(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
{
     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
{
     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
{
     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)));
}

fn body_velocity_at_point_0( com_1 : vec3<f32>,  vels_0 : ptr<function, BodyVelocity_std430_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 : GridHashMapEntryGeneric_std430_0 = entryPointParams_hmap_entries_0[slot_0];
        if((_S16.state_0) == _S15)
        {
            var _S17 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
            return _S17;
        }
        else
        {
            if((_S16.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 fetch_max_linked_lists_length_0( _S18 : vec3<u32>,  _S19 : ptr<function, BlockVirtualId_std430_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 : ptr<function, BlockVirtualId_std430_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;
                var _S46 : bool = (shared_nodes_0[_S44].particle_id_0) != u32(4294967295);
                var _S47 : bool;
                if(_S46)
                {
                    _S47 = (entryPointParams_particles_dyn_0[_S45].enabled_0) != u32(0);
                }
                else
                {
                    _S47 = false;
                }
                if(_S47)
                {
                    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);
                }
                else
                {
                    shared_affinities_0[_S44] = u32(0);
                    var _S48 : vec3<f32> = vec3<f32>(0.0f);
                    shared_normals_0[_S44] = _S48;
                    shared_pos_0[_S44].pt_0 = _S48;
                    shared_affine_0[_S44] = mat3x3<f32>(_S48, _S48, _S48);
                    shared_vel_mass_0[_S44] = vec4<f32>(0.0f);
                }
                if(_S46)
                {
                    shared_nodes_0[_S44].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S45];
                }
                k_2 = k_2 + u32(1);
            }
            j_2 = j_2 + u32(1);
        }
        i_6 = i_6 + u32(1);
    }
    return;
}

fn p2g_step_0( _S49 : u32,  _S50 : f32,  _S51 : u32,  _S52 : u32) -> P2GStepResult_0
{
    var impulse_2 : vec3<f32>;
    var ang_impulse_2 : vec3<f32>;
    var new_momentum_velocity_mass_2 : vec4<f32>;
    var _S53 : u32;
    var _S54 : vec3<u32>;
    var _S55 : vec3<f32>;
    var _S56 : u32 = flatten_shared_shift_0(u32(2), u32(2), u32(2));
    var _S57 : vec4<f32> = vec4<f32>(0.0f);
    var _S58 : vec3<f32> = vec3<f32>(0.0f);
    for(;;)
    {
        for(;;)
        {
            for(;;)
            {
                var _S59 : u32 = _S49 - _S56;
                _S53 = _S59;
                var _S60 : u32 = _S59 + u32(86);
                var _S61 : Position_0 = shared_pos_0[_S60];
                var _S62 : mat3x3<f32> = shared_affine_0[_S60];
                var _S63 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S60], _S50);
                var _S64 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S63, _S50);
                var _S65 : vec3<f32> = shared_vel_mass_0[_S60].xyz;
                var _S66 : f32 = shared_vel_mass_0[_S60].w;
                const _S67 : vec3<u32> = vec3<u32>(u32(2), u32(2), u32(2));
                _S54 = _S67;
                var _S68 : vec3<u32> = _S67 - vec3<u32>(u32(2), u32(2), u32(2));
                var _S69 : vec3<f32> = _S65 * vec3<f32>(_S66);
                var _S70 : vec3<f32> = vec3<f32>(_S68);
                var _S71 : vec3<f32> = vec3<f32>(_S50);
                _S55 = _S71;
                var _S72 : vec3<f32> = _S63 + _S70 * _S71;
                var _S73 : f32 = _S64[i32(0)][_S68.x] * _S64[i32(1)][_S68.y] * _S64[i32(2)][_S68.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S60]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S74 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S75 : vec3<f32> = _S72 + _S61.pt_0;
                        var _S76 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S74), _S75);
                        var _S77 : vec3<f32> = (_S65 - (_S76 + project_velocity_0(_S65 - _S76, shared_normals_0[_S60]))) * vec3<f32>((_S73 * _S66));
                        var _S78 : vec3<f32> = cross(_S77, entryPointParams_body_impulses_0[_S52].com_0 - _S75);
                        new_momentum_velocity_mass_2 = _S57;
                        ang_impulse_2 = _S78;
                        impulse_2 = _S77;
                        break;
                    }
                    new_momentum_velocity_mass_2 = _S57;
                }
                else
                {
                    new_momentum_velocity_mass_2 = vec4<f32>((((_S62) * (_S72))) + _S69, _S66) * vec4<f32>(_S73);
                }
                ang_impulse_2 = _S58;
                impulse_2 = _S58;
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S79 : u32 = _S53 + u32(74);
                var _S80 : Position_0 = shared_pos_0[_S79];
                var _S81 : mat3x3<f32> = shared_affine_0[_S79];
                var _S82 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S79], _S50);
                var _S83 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S82, _S50);
                var _S84 : vec3<f32> = shared_vel_mass_0[_S79].xyz;
                var _S85 : f32 = shared_vel_mass_0[_S79].w;
                var _S86 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(0), u32(2));
                var _S87 : vec3<f32> = _S84 * vec3<f32>(_S85);
                var _S88 : vec3<f32> = _S82 + vec3<f32>(_S86) * _S55;
                var _S89 : f32 = _S83[i32(0)][_S86.x] * _S83[i32(1)][_S86.y] * _S83[i32(2)][_S86.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S79]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S90 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S91 : vec3<f32> = _S88 + _S80.pt_0;
                        var _S92 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S90), _S91);
                        var _S93 : vec3<f32> = (_S84 - (_S92 + project_velocity_0(_S84 - _S92, shared_normals_0[_S79]))) * vec3<f32>((_S89 * _S85));
                        var impulse_3 : vec3<f32> = impulse_2 + _S93;
                        ang_impulse_2 = ang_impulse_2 + cross(_S93, entryPointParams_body_impulses_0[_S52].com_0 - _S91);
                        impulse_2 = impulse_3;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S81) * (_S88))) + _S87, _S85) * vec4<f32>(_S89);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S94 : u32 = _S53 + u32(80);
                var _S95 : Position_0 = shared_pos_0[_S94];
                var _S96 : mat3x3<f32> = shared_affine_0[_S94];
                var _S97 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S94], _S50);
                var _S98 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S97, _S50);
                var _S99 : vec3<f32> = shared_vel_mass_0[_S94].xyz;
                var _S100 : f32 = shared_vel_mass_0[_S94].w;
                var _S101 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(1), u32(2));
                var _S102 : vec3<f32> = _S99 * vec3<f32>(_S100);
                var _S103 : vec3<f32> = _S97 + vec3<f32>(_S101) * _S55;
                var _S104 : f32 = _S98[i32(0)][_S101.x] * _S98[i32(1)][_S101.y] * _S98[i32(2)][_S101.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S94]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S105 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S106 : vec3<f32> = _S103 + _S95.pt_0;
                        var _S107 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S105), _S106);
                        var _S108 : vec3<f32> = (_S99 - (_S107 + project_velocity_0(_S99 - _S107, shared_normals_0[_S94]))) * vec3<f32>((_S104 * _S100));
                        var impulse_4 : vec3<f32> = impulse_2 + _S108;
                        ang_impulse_2 = ang_impulse_2 + cross(_S108, entryPointParams_body_impulses_0[_S52].com_0 - _S106);
                        impulse_2 = impulse_4;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S96) * (_S103))) + _S102, _S100) * vec4<f32>(_S104);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S109 : u32 = _S53 + u32(84);
                var _S110 : Position_0 = shared_pos_0[_S109];
                var _S111 : mat3x3<f32> = shared_affine_0[_S109];
                var _S112 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S109], _S50);
                var _S113 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S112, _S50);
                var _S114 : vec3<f32> = shared_vel_mass_0[_S109].xyz;
                var _S115 : f32 = shared_vel_mass_0[_S109].w;
                var _S116 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(2), u32(2));
                var _S117 : vec3<f32> = _S114 * vec3<f32>(_S115);
                var _S118 : vec3<f32> = _S112 + vec3<f32>(_S116) * _S55;
                var _S119 : f32 = _S113[i32(0)][_S116.x] * _S113[i32(1)][_S116.y] * _S113[i32(2)][_S116.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S109]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S120 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S121 : vec3<f32> = _S118 + _S110.pt_0;
                        var _S122 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S120), _S121);
                        var _S123 : vec3<f32> = (_S114 - (_S122 + project_velocity_0(_S114 - _S122, shared_normals_0[_S109]))) * vec3<f32>((_S119 * _S115));
                        var impulse_5 : vec3<f32> = impulse_2 + _S123;
                        ang_impulse_2 = ang_impulse_2 + cross(_S123, entryPointParams_body_impulses_0[_S52].com_0 - _S121);
                        impulse_2 = impulse_5;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S111) * (_S118))) + _S117, _S115) * vec4<f32>(_S119);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S124 : u32 = _S53 + u32(72);
                var _S125 : Position_0 = shared_pos_0[_S124];
                var _S126 : mat3x3<f32> = shared_affine_0[_S124];
                var _S127 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S124], _S50);
                var _S128 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S127, _S50);
                var _S129 : vec3<f32> = shared_vel_mass_0[_S124].xyz;
                var _S130 : f32 = shared_vel_mass_0[_S124].w;
                var _S131 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(0), u32(2));
                var _S132 : vec3<f32> = _S129 * vec3<f32>(_S130);
                var _S133 : vec3<f32> = _S127 + vec3<f32>(_S131) * _S55;
                var _S134 : f32 = _S128[i32(0)][_S131.x] * _S128[i32(1)][_S131.y] * _S128[i32(2)][_S131.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S124]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S135 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S136 : vec3<f32> = _S133 + _S125.pt_0;
                        var _S137 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S135), _S136);
                        var _S138 : vec3<f32> = (_S129 - (_S137 + project_velocity_0(_S129 - _S137, shared_normals_0[_S124]))) * vec3<f32>((_S134 * _S130));
                        var impulse_6 : vec3<f32> = impulse_2 + _S138;
                        ang_impulse_2 = ang_impulse_2 + cross(_S138, entryPointParams_body_impulses_0[_S52].com_0 - _S136);
                        impulse_2 = impulse_6;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S126) * (_S133))) + _S132, _S130) * vec4<f32>(_S134);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S139 : u32 = _S53 + u32(78);
                var _S140 : Position_0 = shared_pos_0[_S139];
                var _S141 : mat3x3<f32> = shared_affine_0[_S139];
                var _S142 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S139], _S50);
                var _S143 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S142, _S50);
                var _S144 : vec3<f32> = shared_vel_mass_0[_S139].xyz;
                var _S145 : f32 = shared_vel_mass_0[_S139].w;
                var _S146 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(1), u32(2));
                var _S147 : vec3<f32> = _S144 * vec3<f32>(_S145);
                var _S148 : vec3<f32> = _S142 + vec3<f32>(_S146) * _S55;
                var _S149 : f32 = _S143[i32(0)][_S146.x] * _S143[i32(1)][_S146.y] * _S143[i32(2)][_S146.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S139]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S150 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S151 : vec3<f32> = _S148 + _S140.pt_0;
                        var _S152 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S150), _S151);
                        var _S153 : vec3<f32> = (_S144 - (_S152 + project_velocity_0(_S144 - _S152, shared_normals_0[_S139]))) * vec3<f32>((_S149 * _S145));
                        var impulse_7 : vec3<f32> = impulse_2 + _S153;
                        ang_impulse_2 = ang_impulse_2 + cross(_S153, entryPointParams_body_impulses_0[_S52].com_0 - _S151);
                        impulse_2 = impulse_7;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S141) * (_S148))) + _S147, _S145) * vec4<f32>(_S149);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S154 : u32 = _S53 + u32(85);
                var _S155 : Position_0 = shared_pos_0[_S154];
                var _S156 : mat3x3<f32> = shared_affine_0[_S154];
                var _S157 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S154], _S50);
                var _S158 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S157, _S50);
                var _S159 : vec3<f32> = shared_vel_mass_0[_S154].xyz;
                var _S160 : f32 = shared_vel_mass_0[_S154].w;
                var _S161 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(2), u32(2));
                var _S162 : vec3<f32> = _S159 * vec3<f32>(_S160);
                var _S163 : vec3<f32> = _S157 + vec3<f32>(_S161) * _S55;
                var _S164 : f32 = _S158[i32(0)][_S161.x] * _S158[i32(1)][_S161.y] * _S158[i32(2)][_S161.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S154]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S165 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S166 : vec3<f32> = _S163 + _S155.pt_0;
                        var _S167 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S165), _S166);
                        var _S168 : vec3<f32> = (_S159 - (_S167 + project_velocity_0(_S159 - _S167, shared_normals_0[_S154]))) * vec3<f32>((_S164 * _S160));
                        var impulse_8 : vec3<f32> = impulse_2 + _S168;
                        ang_impulse_2 = ang_impulse_2 + cross(_S168, entryPointParams_body_impulses_0[_S52].com_0 - _S166);
                        impulse_2 = impulse_8;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S156) * (_S163))) + _S162, _S160) * vec4<f32>(_S164);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S169 : u32 = _S53 + u32(73);
                var _S170 : Position_0 = shared_pos_0[_S169];
                var _S171 : mat3x3<f32> = shared_affine_0[_S169];
                var _S172 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S169], _S50);
                var _S173 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S172, _S50);
                var _S174 : vec3<f32> = shared_vel_mass_0[_S169].xyz;
                var _S175 : f32 = shared_vel_mass_0[_S169].w;
                var _S176 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(0), u32(2));
                var _S177 : vec3<f32> = _S174 * vec3<f32>(_S175);
                var _S178 : vec3<f32> = _S172 + vec3<f32>(_S176) * _S55;
                var _S179 : f32 = _S173[i32(0)][_S176.x] * _S173[i32(1)][_S176.y] * _S173[i32(2)][_S176.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S169]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S180 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S181 : vec3<f32> = _S178 + _S170.pt_0;
                        var _S182 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S180), _S181);
                        var _S183 : vec3<f32> = (_S174 - (_S182 + project_velocity_0(_S174 - _S182, shared_normals_0[_S169]))) * vec3<f32>((_S179 * _S175));
                        var impulse_9 : vec3<f32> = impulse_2 + _S183;
                        ang_impulse_2 = ang_impulse_2 + cross(_S183, entryPointParams_body_impulses_0[_S52].com_0 - _S181);
                        impulse_2 = impulse_9;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S171) * (_S178))) + _S177, _S175) * vec4<f32>(_S179);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S184 : u32 = _S53 + u32(79);
                var _S185 : Position_0 = shared_pos_0[_S184];
                var _S186 : mat3x3<f32> = shared_affine_0[_S184];
                var _S187 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S184], _S50);
                var _S188 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S187, _S50);
                var _S189 : vec3<f32> = shared_vel_mass_0[_S184].xyz;
                var _S190 : f32 = shared_vel_mass_0[_S184].w;
                var _S191 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(1), u32(2));
                var _S192 : vec3<f32> = _S189 * vec3<f32>(_S190);
                var _S193 : vec3<f32> = _S187 + vec3<f32>(_S191) * _S55;
                var _S194 : f32 = _S188[i32(0)][_S191.x] * _S188[i32(1)][_S191.y] * _S188[i32(2)][_S191.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S184]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S195 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S196 : vec3<f32> = _S193 + _S185.pt_0;
                        var _S197 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S195), _S196);
                        var _S198 : vec3<f32> = (_S189 - (_S197 + project_velocity_0(_S189 - _S197, shared_normals_0[_S184]))) * vec3<f32>((_S194 * _S190));
                        var impulse_10 : vec3<f32> = impulse_2 + _S198;
                        ang_impulse_2 = ang_impulse_2 + cross(_S198, entryPointParams_body_impulses_0[_S52].com_0 - _S196);
                        impulse_2 = impulse_10;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S186) * (_S193))) + _S192, _S190) * vec4<f32>(_S194);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S199 : u32 = _S53 + u32(14);
                var _S200 : Position_0 = shared_pos_0[_S199];
                var _S201 : mat3x3<f32> = shared_affine_0[_S199];
                var _S202 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S199], _S50);
                var _S203 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S202, _S50);
                var _S204 : vec3<f32> = shared_vel_mass_0[_S199].xyz;
                var _S205 : f32 = shared_vel_mass_0[_S199].w;
                var _S206 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(2), u32(0));
                var _S207 : vec3<f32> = _S204 * vec3<f32>(_S205);
                var _S208 : vec3<f32> = _S202 + vec3<f32>(_S206) * _S55;
                var _S209 : f32 = _S203[i32(0)][_S206.x] * _S203[i32(1)][_S206.y] * _S203[i32(2)][_S206.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S199]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S210 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S211 : vec3<f32> = _S208 + _S200.pt_0;
                        var _S212 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S210), _S211);
                        var _S213 : vec3<f32> = (_S204 - (_S212 + project_velocity_0(_S204 - _S212, shared_normals_0[_S199]))) * vec3<f32>((_S209 * _S205));
                        var impulse_11 : vec3<f32> = impulse_2 + _S213;
                        ang_impulse_2 = ang_impulse_2 + cross(_S213, entryPointParams_body_impulses_0[_S52].com_0 - _S211);
                        impulse_2 = impulse_11;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S201) * (_S208))) + _S207, _S205) * vec4<f32>(_S209);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S214 : u32 = _S53 + u32(2);
                var _S215 : Position_0 = shared_pos_0[_S214];
                var _S216 : mat3x3<f32> = shared_affine_0[_S214];
                var _S217 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S214], _S50);
                var _S218 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S217, _S50);
                var _S219 : vec3<f32> = shared_vel_mass_0[_S214].xyz;
                var _S220 : f32 = shared_vel_mass_0[_S214].w;
                var _S221 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(0), u32(0));
                var _S222 : vec3<f32> = _S219 * vec3<f32>(_S220);
                var _S223 : vec3<f32> = _S217 + vec3<f32>(_S221) * _S55;
                var _S224 : f32 = _S218[i32(0)][_S221.x] * _S218[i32(1)][_S221.y] * _S218[i32(2)][_S221.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S214]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S225 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S226 : vec3<f32> = _S223 + _S215.pt_0;
                        var _S227 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S225), _S226);
                        var _S228 : vec3<f32> = (_S219 - (_S227 + project_velocity_0(_S219 - _S227, shared_normals_0[_S214]))) * vec3<f32>((_S224 * _S220));
                        var impulse_12 : vec3<f32> = impulse_2 + _S228;
                        ang_impulse_2 = ang_impulse_2 + cross(_S228, entryPointParams_body_impulses_0[_S52].com_0 - _S226);
                        impulse_2 = impulse_12;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S216) * (_S223))) + _S222, _S220) * vec4<f32>(_S224);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S229 : u32 = _S53 + u32(8);
                var _S230 : Position_0 = shared_pos_0[_S229];
                var _S231 : mat3x3<f32> = shared_affine_0[_S229];
                var _S232 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S229], _S50);
                var _S233 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S232, _S50);
                var _S234 : vec3<f32> = shared_vel_mass_0[_S229].xyz;
                var _S235 : f32 = shared_vel_mass_0[_S229].w;
                var _S236 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(1), u32(0));
                var _S237 : vec3<f32> = _S234 * vec3<f32>(_S235);
                var _S238 : vec3<f32> = _S232 + vec3<f32>(_S236) * _S55;
                var _S239 : f32 = _S233[i32(0)][_S236.x] * _S233[i32(1)][_S236.y] * _S233[i32(2)][_S236.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S229]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S240 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S241 : vec3<f32> = _S238 + _S230.pt_0;
                        var _S242 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S240), _S241);
                        var _S243 : vec3<f32> = (_S234 - (_S242 + project_velocity_0(_S234 - _S242, shared_normals_0[_S229]))) * vec3<f32>((_S239 * _S235));
                        var impulse_13 : vec3<f32> = impulse_2 + _S243;
                        ang_impulse_2 = ang_impulse_2 + cross(_S243, entryPointParams_body_impulses_0[_S52].com_0 - _S241);
                        impulse_2 = impulse_13;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S231) * (_S238))) + _S237, _S235) * vec4<f32>(_S239);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S244 : u32 = _S53 + u32(12);
                var _S245 : Position_0 = shared_pos_0[_S244];
                var _S246 : mat3x3<f32> = shared_affine_0[_S244];
                var _S247 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S244], _S50);
                var _S248 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S247, _S50);
                var _S249 : vec3<f32> = shared_vel_mass_0[_S244].xyz;
                var _S250 : f32 = shared_vel_mass_0[_S244].w;
                var _S251 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(2), u32(0));
                var _S252 : vec3<f32> = _S249 * vec3<f32>(_S250);
                var _S253 : vec3<f32> = _S247 + vec3<f32>(_S251) * _S55;
                var _S254 : f32 = _S248[i32(0)][_S251.x] * _S248[i32(1)][_S251.y] * _S248[i32(2)][_S251.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S244]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S255 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S256 : vec3<f32> = _S253 + _S245.pt_0;
                        var _S257 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S255), _S256);
                        var _S258 : vec3<f32> = (_S249 - (_S257 + project_velocity_0(_S249 - _S257, shared_normals_0[_S244]))) * vec3<f32>((_S254 * _S250));
                        var impulse_14 : vec3<f32> = impulse_2 + _S258;
                        ang_impulse_2 = ang_impulse_2 + cross(_S258, entryPointParams_body_impulses_0[_S52].com_0 - _S256);
                        impulse_2 = impulse_14;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S246) * (_S253))) + _S252, _S250) * vec4<f32>(_S254);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S259 : Position_0 = shared_pos_0[_S53];
                var _S260 : mat3x3<f32> = shared_affine_0[_S53];
                var _S261 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S53], _S50);
                var _S262 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S261, _S50);
                var _S263 : vec3<f32> = shared_vel_mass_0[_S53].xyz;
                var _S264 : f32 = shared_vel_mass_0[_S53].w;
                var _S265 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(0), u32(0));
                var _S266 : vec3<f32> = _S263 * vec3<f32>(_S264);
                var _S267 : vec3<f32> = _S261 + vec3<f32>(_S265) * _S55;
                var _S268 : f32 = _S262[i32(0)][_S265.x] * _S262[i32(1)][_S265.y] * _S262[i32(2)][_S265.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S53]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S269 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S270 : vec3<f32> = _S267 + _S259.pt_0;
                        var _S271 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S269), _S270);
                        var _S272 : vec3<f32> = (_S263 - (_S271 + project_velocity_0(_S263 - _S271, shared_normals_0[_S53]))) * vec3<f32>((_S268 * _S264));
                        var impulse_15 : vec3<f32> = impulse_2 + _S272;
                        ang_impulse_2 = ang_impulse_2 + cross(_S272, entryPointParams_body_impulses_0[_S52].com_0 - _S270);
                        impulse_2 = impulse_15;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S260) * (_S267))) + _S266, _S264) * vec4<f32>(_S268);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S273 : u32 = _S53 + u32(6);
                var _S274 : Position_0 = shared_pos_0[_S273];
                var _S275 : mat3x3<f32> = shared_affine_0[_S273];
                var _S276 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S273], _S50);
                var _S277 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S276, _S50);
                var _S278 : vec3<f32> = shared_vel_mass_0[_S273].xyz;
                var _S279 : f32 = shared_vel_mass_0[_S273].w;
                var _S280 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(1), u32(0));
                var _S281 : vec3<f32> = _S278 * vec3<f32>(_S279);
                var _S282 : vec3<f32> = _S276 + vec3<f32>(_S280) * _S55;
                var _S283 : f32 = _S277[i32(0)][_S280.x] * _S277[i32(1)][_S280.y] * _S277[i32(2)][_S280.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S273]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S284 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S285 : vec3<f32> = _S282 + _S274.pt_0;
                        var _S286 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S284), _S285);
                        var _S287 : vec3<f32> = (_S278 - (_S286 + project_velocity_0(_S278 - _S286, shared_normals_0[_S273]))) * vec3<f32>((_S283 * _S279));
                        var impulse_16 : vec3<f32> = impulse_2 + _S287;
                        ang_impulse_2 = ang_impulse_2 + cross(_S287, entryPointParams_body_impulses_0[_S52].com_0 - _S285);
                        impulse_2 = impulse_16;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S275) * (_S282))) + _S281, _S279) * vec4<f32>(_S283);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S288 : u32 = _S53 + u32(13);
                var _S289 : Position_0 = shared_pos_0[_S288];
                var _S290 : mat3x3<f32> = shared_affine_0[_S288];
                var _S291 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S288], _S50);
                var _S292 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S291, _S50);
                var _S293 : vec3<f32> = shared_vel_mass_0[_S288].xyz;
                var _S294 : f32 = shared_vel_mass_0[_S288].w;
                var _S295 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(2), u32(0));
                var _S296 : vec3<f32> = _S293 * vec3<f32>(_S294);
                var _S297 : vec3<f32> = _S291 + vec3<f32>(_S295) * _S55;
                var _S298 : f32 = _S292[i32(0)][_S295.x] * _S292[i32(1)][_S295.y] * _S292[i32(2)][_S295.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S288]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S299 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S300 : vec3<f32> = _S297 + _S289.pt_0;
                        var _S301 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S299), _S300);
                        var _S302 : vec3<f32> = (_S293 - (_S301 + project_velocity_0(_S293 - _S301, shared_normals_0[_S288]))) * vec3<f32>((_S298 * _S294));
                        var impulse_17 : vec3<f32> = impulse_2 + _S302;
                        ang_impulse_2 = ang_impulse_2 + cross(_S302, entryPointParams_body_impulses_0[_S52].com_0 - _S300);
                        impulse_2 = impulse_17;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S290) * (_S297))) + _S296, _S294) * vec4<f32>(_S298);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S303 : u32 = _S53 + u32(1);
                var _S304 : Position_0 = shared_pos_0[_S303];
                var _S305 : mat3x3<f32> = shared_affine_0[_S303];
                var _S306 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S303], _S50);
                var _S307 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S306, _S50);
                var _S308 : vec3<f32> = shared_vel_mass_0[_S303].xyz;
                var _S309 : f32 = shared_vel_mass_0[_S303].w;
                var _S310 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(0), u32(0));
                var _S311 : vec3<f32> = _S308 * vec3<f32>(_S309);
                var _S312 : vec3<f32> = _S306 + vec3<f32>(_S310) * _S55;
                var _S313 : f32 = _S307[i32(0)][_S310.x] * _S307[i32(1)][_S310.y] * _S307[i32(2)][_S310.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S303]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S314 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S315 : vec3<f32> = _S312 + _S304.pt_0;
                        var _S316 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S314), _S315);
                        var _S317 : vec3<f32> = (_S308 - (_S316 + project_velocity_0(_S308 - _S316, shared_normals_0[_S303]))) * vec3<f32>((_S313 * _S309));
                        var impulse_18 : vec3<f32> = impulse_2 + _S317;
                        ang_impulse_2 = ang_impulse_2 + cross(_S317, entryPointParams_body_impulses_0[_S52].com_0 - _S315);
                        impulse_2 = impulse_18;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S305) * (_S312))) + _S311, _S309) * vec4<f32>(_S313);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S318 : u32 = _S53 + u32(7);
                var _S319 : Position_0 = shared_pos_0[_S318];
                var _S320 : mat3x3<f32> = shared_affine_0[_S318];
                var _S321 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S318], _S50);
                var _S322 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S321, _S50);
                var _S323 : vec3<f32> = shared_vel_mass_0[_S318].xyz;
                var _S324 : f32 = shared_vel_mass_0[_S318].w;
                var _S325 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(1), u32(0));
                var _S326 : vec3<f32> = _S323 * vec3<f32>(_S324);
                var _S327 : vec3<f32> = _S321 + vec3<f32>(_S325) * _S55;
                var _S328 : f32 = _S322[i32(0)][_S325.x] * _S322[i32(1)][_S325.y] * _S322[i32(2)][_S325.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S318]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S329 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S330 : vec3<f32> = _S327 + _S319.pt_0;
                        var _S331 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S329), _S330);
                        var _S332 : vec3<f32> = (_S323 - (_S331 + project_velocity_0(_S323 - _S331, shared_normals_0[_S318]))) * vec3<f32>((_S328 * _S324));
                        var impulse_19 : vec3<f32> = impulse_2 + _S332;
                        ang_impulse_2 = ang_impulse_2 + cross(_S332, entryPointParams_body_impulses_0[_S52].com_0 - _S330);
                        impulse_2 = impulse_19;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S320) * (_S327))) + _S326, _S324) * vec4<f32>(_S328);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S333 : u32 = _S53 + u32(50);
                var _S334 : Position_0 = shared_pos_0[_S333];
                var _S335 : mat3x3<f32> = shared_affine_0[_S333];
                var _S336 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S333], _S50);
                var _S337 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S336, _S50);
                var _S338 : vec3<f32> = shared_vel_mass_0[_S333].xyz;
                var _S339 : f32 = shared_vel_mass_0[_S333].w;
                var _S340 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(2), u32(1));
                var _S341 : vec3<f32> = _S338 * vec3<f32>(_S339);
                var _S342 : vec3<f32> = _S336 + vec3<f32>(_S340) * _S55;
                var _S343 : f32 = _S337[i32(0)][_S340.x] * _S337[i32(1)][_S340.y] * _S337[i32(2)][_S340.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S333]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S344 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S345 : vec3<f32> = _S342 + _S334.pt_0;
                        var _S346 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S344), _S345);
                        var _S347 : vec3<f32> = (_S338 - (_S346 + project_velocity_0(_S338 - _S346, shared_normals_0[_S333]))) * vec3<f32>((_S343 * _S339));
                        var impulse_20 : vec3<f32> = impulse_2 + _S347;
                        ang_impulse_2 = ang_impulse_2 + cross(_S347, entryPointParams_body_impulses_0[_S52].com_0 - _S345);
                        impulse_2 = impulse_20;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S335) * (_S342))) + _S341, _S339) * vec4<f32>(_S343);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S348 : u32 = _S53 + u32(38);
                var _S349 : Position_0 = shared_pos_0[_S348];
                var _S350 : mat3x3<f32> = shared_affine_0[_S348];
                var _S351 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S348], _S50);
                var _S352 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S351, _S50);
                var _S353 : vec3<f32> = shared_vel_mass_0[_S348].xyz;
                var _S354 : f32 = shared_vel_mass_0[_S348].w;
                var _S355 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(0), u32(1));
                var _S356 : vec3<f32> = _S353 * vec3<f32>(_S354);
                var _S357 : vec3<f32> = _S351 + vec3<f32>(_S355) * _S55;
                var _S358 : f32 = _S352[i32(0)][_S355.x] * _S352[i32(1)][_S355.y] * _S352[i32(2)][_S355.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S348]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S359 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S360 : vec3<f32> = _S357 + _S349.pt_0;
                        var _S361 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S359), _S360);
                        var _S362 : vec3<f32> = (_S353 - (_S361 + project_velocity_0(_S353 - _S361, shared_normals_0[_S348]))) * vec3<f32>((_S358 * _S354));
                        var impulse_21 : vec3<f32> = impulse_2 + _S362;
                        ang_impulse_2 = ang_impulse_2 + cross(_S362, entryPointParams_body_impulses_0[_S52].com_0 - _S360);
                        impulse_2 = impulse_21;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S350) * (_S357))) + _S356, _S354) * vec4<f32>(_S358);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S363 : u32 = _S53 + u32(44);
                var _S364 : Position_0 = shared_pos_0[_S363];
                var _S365 : mat3x3<f32> = shared_affine_0[_S363];
                var _S366 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S363], _S50);
                var _S367 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S366, _S50);
                var _S368 : vec3<f32> = shared_vel_mass_0[_S363].xyz;
                var _S369 : f32 = shared_vel_mass_0[_S363].w;
                var _S370 : vec3<u32> = _S54 - vec3<u32>(u32(2), u32(1), u32(1));
                var _S371 : vec3<f32> = _S368 * vec3<f32>(_S369);
                var _S372 : vec3<f32> = _S366 + vec3<f32>(_S370) * _S55;
                var _S373 : f32 = _S367[i32(0)][_S370.x] * _S367[i32(1)][_S370.y] * _S367[i32(2)][_S370.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S363]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S374 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S375 : vec3<f32> = _S372 + _S364.pt_0;
                        var _S376 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S374), _S375);
                        var _S377 : vec3<f32> = (_S368 - (_S376 + project_velocity_0(_S368 - _S376, shared_normals_0[_S363]))) * vec3<f32>((_S373 * _S369));
                        var impulse_22 : vec3<f32> = impulse_2 + _S377;
                        ang_impulse_2 = ang_impulse_2 + cross(_S377, entryPointParams_body_impulses_0[_S52].com_0 - _S375);
                        impulse_2 = impulse_22;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S365) * (_S372))) + _S371, _S369) * vec4<f32>(_S373);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S378 : u32 = _S53 + u32(48);
                var _S379 : Position_0 = shared_pos_0[_S378];
                var _S380 : mat3x3<f32> = shared_affine_0[_S378];
                var _S381 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S378], _S50);
                var _S382 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S381, _S50);
                var _S383 : vec3<f32> = shared_vel_mass_0[_S378].xyz;
                var _S384 : f32 = shared_vel_mass_0[_S378].w;
                var _S385 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(2), u32(1));
                var _S386 : vec3<f32> = _S383 * vec3<f32>(_S384);
                var _S387 : vec3<f32> = _S381 + vec3<f32>(_S385) * _S55;
                var _S388 : f32 = _S382[i32(0)][_S385.x] * _S382[i32(1)][_S385.y] * _S382[i32(2)][_S385.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S378]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S389 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S390 : vec3<f32> = _S387 + _S379.pt_0;
                        var _S391 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S389), _S390);
                        var _S392 : vec3<f32> = (_S383 - (_S391 + project_velocity_0(_S383 - _S391, shared_normals_0[_S378]))) * vec3<f32>((_S388 * _S384));
                        var impulse_23 : vec3<f32> = impulse_2 + _S392;
                        ang_impulse_2 = ang_impulse_2 + cross(_S392, entryPointParams_body_impulses_0[_S52].com_0 - _S390);
                        impulse_2 = impulse_23;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S380) * (_S387))) + _S386, _S384) * vec4<f32>(_S388);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S393 : u32 = _S53 + u32(36);
                var _S394 : Position_0 = shared_pos_0[_S393];
                var _S395 : mat3x3<f32> = shared_affine_0[_S393];
                var _S396 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S393], _S50);
                var _S397 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S396, _S50);
                var _S398 : vec3<f32> = shared_vel_mass_0[_S393].xyz;
                var _S399 : f32 = shared_vel_mass_0[_S393].w;
                var _S400 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(0), u32(1));
                var _S401 : vec3<f32> = _S398 * vec3<f32>(_S399);
                var _S402 : vec3<f32> = _S396 + vec3<f32>(_S400) * _S55;
                var _S403 : f32 = _S397[i32(0)][_S400.x] * _S397[i32(1)][_S400.y] * _S397[i32(2)][_S400.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S393]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S404 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S405 : vec3<f32> = _S402 + _S394.pt_0;
                        var _S406 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S404), _S405);
                        var _S407 : vec3<f32> = (_S398 - (_S406 + project_velocity_0(_S398 - _S406, shared_normals_0[_S393]))) * vec3<f32>((_S403 * _S399));
                        var impulse_24 : vec3<f32> = impulse_2 + _S407;
                        ang_impulse_2 = ang_impulse_2 + cross(_S407, entryPointParams_body_impulses_0[_S52].com_0 - _S405);
                        impulse_2 = impulse_24;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S395) * (_S402))) + _S401, _S399) * vec4<f32>(_S403);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S408 : u32 = _S53 + u32(42);
                var _S409 : Position_0 = shared_pos_0[_S408];
                var _S410 : mat3x3<f32> = shared_affine_0[_S408];
                var _S411 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S408], _S50);
                var _S412 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S411, _S50);
                var _S413 : vec3<f32> = shared_vel_mass_0[_S408].xyz;
                var _S414 : f32 = shared_vel_mass_0[_S408].w;
                var _S415 : vec3<u32> = _S54 - vec3<u32>(u32(0), u32(1), u32(1));
                var _S416 : vec3<f32> = _S413 * vec3<f32>(_S414);
                var _S417 : vec3<f32> = _S411 + vec3<f32>(_S415) * _S55;
                var _S418 : f32 = _S412[i32(0)][_S415.x] * _S412[i32(1)][_S415.y] * _S412[i32(2)][_S415.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S408]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S419 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S420 : vec3<f32> = _S417 + _S409.pt_0;
                        var _S421 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S419), _S420);
                        var _S422 : vec3<f32> = (_S413 - (_S421 + project_velocity_0(_S413 - _S421, shared_normals_0[_S408]))) * vec3<f32>((_S418 * _S414));
                        var impulse_25 : vec3<f32> = impulse_2 + _S422;
                        ang_impulse_2 = ang_impulse_2 + cross(_S422, entryPointParams_body_impulses_0[_S52].com_0 - _S420);
                        impulse_2 = impulse_25;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S410) * (_S417))) + _S416, _S414) * vec4<f32>(_S418);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S423 : u32 = _S53 + u32(49);
                var _S424 : Position_0 = shared_pos_0[_S423];
                var _S425 : mat3x3<f32> = shared_affine_0[_S423];
                var _S426 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S423], _S50);
                var _S427 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S426, _S50);
                var _S428 : vec3<f32> = shared_vel_mass_0[_S423].xyz;
                var _S429 : f32 = shared_vel_mass_0[_S423].w;
                var _S430 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(2), u32(1));
                var _S431 : vec3<f32> = _S428 * vec3<f32>(_S429);
                var _S432 : vec3<f32> = _S426 + vec3<f32>(_S430) * _S55;
                var _S433 : f32 = _S427[i32(0)][_S430.x] * _S427[i32(1)][_S430.y] * _S427[i32(2)][_S430.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S423]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S434 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S435 : vec3<f32> = _S432 + _S424.pt_0;
                        var _S436 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S434), _S435);
                        var _S437 : vec3<f32> = (_S428 - (_S436 + project_velocity_0(_S428 - _S436, shared_normals_0[_S423]))) * vec3<f32>((_S433 * _S429));
                        var impulse_26 : vec3<f32> = impulse_2 + _S437;
                        ang_impulse_2 = ang_impulse_2 + cross(_S437, entryPointParams_body_impulses_0[_S52].com_0 - _S435);
                        impulse_2 = impulse_26;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S425) * (_S432))) + _S431, _S429) * vec4<f32>(_S433);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S438 : u32 = _S53 + u32(37);
                var _S439 : Position_0 = shared_pos_0[_S438];
                var _S440 : mat3x3<f32> = shared_affine_0[_S438];
                var _S441 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S438], _S50);
                var _S442 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S441, _S50);
                var _S443 : vec3<f32> = shared_vel_mass_0[_S438].xyz;
                var _S444 : f32 = shared_vel_mass_0[_S438].w;
                var _S445 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(0), u32(1));
                var _S446 : vec3<f32> = _S443 * vec3<f32>(_S444);
                var _S447 : vec3<f32> = _S441 + vec3<f32>(_S445) * _S55;
                var _S448 : f32 = _S442[i32(0)][_S445.x] * _S442[i32(1)][_S445.y] * _S442[i32(2)][_S445.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S438]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S449 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S450 : vec3<f32> = _S447 + _S439.pt_0;
                        var _S451 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S449), _S450);
                        var _S452 : vec3<f32> = (_S443 - (_S451 + project_velocity_0(_S443 - _S451, shared_normals_0[_S438]))) * vec3<f32>((_S448 * _S444));
                        var impulse_27 : vec3<f32> = impulse_2 + _S452;
                        ang_impulse_2 = ang_impulse_2 + cross(_S452, entryPointParams_body_impulses_0[_S52].com_0 - _S450);
                        impulse_2 = impulse_27;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S440) * (_S447))) + _S446, _S444) * vec4<f32>(_S448);
                }
                break;
            }
            break;
        }
        for(;;)
        {
            for(;;)
            {
                var _S453 : u32 = _S53 + u32(43);
                var _S454 : Position_0 = shared_pos_0[_S453];
                var _S455 : mat3x3<f32> = shared_affine_0[_S453];
                var _S456 : vec3<f32> = dir_to_associated_grid_node_0(shared_pos_0[_S453], _S50);
                var _S457 : mat3x3<f32> = QuadraticKernel_precompute_weights_0(_S456, _S50);
                var _S458 : vec3<f32> = shared_vel_mass_0[_S453].xyz;
                var _S459 : f32 = shared_vel_mass_0[_S453].w;
                var _S460 : vec3<u32> = _S54 - vec3<u32>(u32(1), u32(1), u32(1));
                var _S461 : vec3<f32> = _S458 * vec3<f32>(_S459);
                var _S462 : vec3<f32> = _S456 + vec3<f32>(_S460) * _S55;
                var _S463 : f32 = _S457[i32(0)][_S460.x] * _S457[i32(1)][_S460.y] * _S457[i32(2)][_S460.z];
                if(!affinities_are_compatible_0(_S51, shared_affinities_0[_S453]))
                {
                    if(_S52 != u32(4294967295))
                    {
                        var _S464 : BodyVelocity_std430_0 = entryPointParams_body_vels_0[_S52];
                        var _S465 : vec3<f32> = _S462 + _S454.pt_0;
                        var _S466 : vec3<f32> = body_velocity_at_point_0(entryPointParams_body_impulses_0[_S52].com_0, &(_S464), _S465);
                        var _S467 : vec3<f32> = (_S458 - (_S466 + project_velocity_0(_S458 - _S466, shared_normals_0[_S453]))) * vec3<f32>((_S463 * _S459));
                        var impulse_28 : vec3<f32> = impulse_2 + _S467;
                        ang_impulse_2 = ang_impulse_2 + cross(_S467, entryPointParams_body_impulses_0[_S52].com_0 - _S465);
                        impulse_2 = impulse_28;
                        break;
                    }
                }
                else
                {
                    new_momentum_velocity_mass_2 = new_momentum_velocity_mass_2 + vec4<f32>((((_S455) * (_S462))) + _S461, _S459) * vec4<f32>(_S463);
                }
                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 _S468 : u32 = block_id_0.x;
    var _S469 : ActiveBlockHeaderGeneric_std430_0 = entryPointParams_active_blocks_0[_S468];
    if(tid_flat_0 == u32(0))
    {
        atomicStore(&(max_linked_list_length_0), u32(0));
    }
    workgroupBarrier();
    fetch_max_linked_lists_length_0(tid_0, &(_S469.virtual_id_0), _S468);
    workgroupBarrier();
    var _S470 : u32 = atomicLoad(&(max_linked_list_length_0));
    max_linked_list_length_uniform_0 = _S470;
    fetch_nodes_0(tid_0, &(_S469.virtual_id_0), _S468);
    var _S471 : u32 = flatten_shared_index_0(tid_0.x + u32(4), tid_0.y + u32(4), tid_0.z + u32(4));
    var _S472 : u32 = shared_nodes_0[_S471].global_id_0;
    var _S473 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S471].global_id_0].cdf_1.affinities_0;
    var _S474 : u32 = entryPointParams_nodes_0[shared_nodes_0[_S471].global_id_0].cdf_1.closest_id_0;
    var _S475 : vec3<f32> = vec3<f32>(0.0f);
    var total_result_0 : P2GStepResult_0 = P2GStepResult_x24init_0(vec4<f32>(0.0f), _S475, _S475);
    var _S476 : u32 = (workgroupUniformLoad(&((max_linked_list_length_uniform_0))));
    var i_7 : u32 = u32(0);
    for(;;)
    {
        if(i_7 < _S476)
        {
        }
        else
        {
            break;
        }
        workgroupBarrier();
        fetch_next_particle_0(tid_0);
        workgroupBarrier();
        var _S477 : P2GStepResult_0 = p2g_step_0(_S471, entryPointParams_grid_0[i32(0)].cell_width_0, _S473, _S474);
        total_result_0.new_momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0 + _S477.new_momentum_velocity_mass_0;
        total_result_0.impulse_0 = total_result_0.impulse_0 + _S477.impulse_0;
        total_result_0.ang_impulse_0 = total_result_0.ang_impulse_0 + _S477.ang_impulse_0;
        i_7 = i_7 + u32(1);
    }
    entryPointParams_nodes_0[_S472].momentum_velocity_mass_0 = total_result_0.new_momentum_velocity_mass_0;
    if(_S474 != u32(4294967295))
    {
        var _S478 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].linear_x_0), flt2int_0(total_result_0.impulse_0.x));
        var _S479 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].linear_y_0), flt2int_0(total_result_0.impulse_0.y));
        var _S480 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].linear_z_0), flt2int_0(total_result_0.impulse_0.z));
        var _S481 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].angular_x_0), flt2int_0(total_result_0.ang_impulse_0.x));
        var _S482 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].angular_y_0), flt2int_0(total_result_0.ang_impulse_0.y));
        var _S483 : i32 = atomicAdd(&(entryPointParams_body_impulses_0[_S474].angular_z_0), flt2int_0(total_result_0.ang_impulse_0.z));
    }
    return;
}