slosh3d 0.2.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) 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;
                if((shared_nodes_0[_S44].particle_id_0) != u32(4294967295))
                {
                    _S46 = (entryPointParams_particles_dyn_0[_S45].enabled_0) != u32(0);
                }
                else
                {
                    _S46 = false;
                }
                if(_S46)
                {
                    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);
                    shared_nodes_0[_S44].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S45];
                }
                else
                {
                    shared_affinities_0[_S44] = u32(0);
                    var _S47 : vec3<f32> = vec3<f32>(0.0f);
                    shared_normals_0[_S44] = _S47;
                    shared_pos_0[_S44].pt_0 = _S47;
                    shared_affine_0[_S44] = mat3x3<f32>(_S47, _S47, _S47);
                    shared_vel_mass_0[_S44] = vec4<f32>(0.0f);
                }
                k_2 = k_2 + u32(1);
            }
            j_2 = j_2 + u32(1);
        }
        i_6 = i_6 + u32(1);
    }
    return;
}

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