struct GridGeneric_std430_0
{
@align(4) num_active_blocks_0 : atomic<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_write> entryPointParams_grid_0 : array<GridGeneric_std430_0>;
struct BlockVirtualId_std430_0
{
@align(8) id_0 : vec2<i32>,
};
struct BlockHeaderId_std430_0
{
@align(4) id_1 : u32,
};
struct GridHashMapEntryGeneric_std430_0
{
@align(8) state_0 : atomic<u32>,
@align(8) key_0 : BlockVirtualId_std430_0,
@align(8) value_0 : BlockHeaderId_std430_0,
};
@binding(1) @group(0) var<storage, read_write> entryPointParams_hmap_entries_0 : array<GridHashMapEntryGeneric_std430_0>;
struct ActiveBlockHeaderGeneric_std430_0
{
@align(8) virtual_id_0 : BlockVirtualId_std430_0,
@align(8) first_particle_0 : u32,
@align(4) num_particles_0 : u32,
};
@binding(2) @group(0) var<storage, read_write> entryPointParams_active_blocks_0 : array<ActiveBlockHeaderGeneric_std430_0>;
struct Position_std430_0
{
@align(8) pt_0 : vec2<f32>,
};
@binding(3) @group(0) var<storage, read> entryPointParams_particles_pos_0 : array<Position_std430_0>;
@binding(4) @group(0) var<uniform> entryPointParams_particles_len_0 : u32;
@binding(5) @group(0) var<storage, read_write> entryPointParams_grid_1 : array<GridGeneric_std430_0>;
@binding(6) @group(0) var<storage, read_write> entryPointParams_hmap_entries_1 : array<GridHashMapEntryGeneric_std430_0>;
@binding(7) @group(0) var<storage, read_write> entryPointParams_active_blocks_1 : array<ActiveBlockHeaderGeneric_std430_0>;
@binding(8) @group(0) var<storage, read> entryPointParams_rigid_particles_pos_0 : array<Position_std430_0>;
@binding(9) @group(0) var<storage, read> entryPointParams_rigid_particle_needs_block_0 : array<u32>;
struct GridGeneric_std430_1
{
@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(10) @group(0) var<storage, read> entryPointParams_grid_2 : array<GridGeneric_std430_1>;
struct GridHashMapEntryGeneric_std430_1
{
@align(8) state_0 : u32,
@align(8) key_0 : BlockVirtualId_std430_0,
@align(8) value_0 : BlockHeaderId_std430_0,
};
@binding(11) @group(0) var<storage, read> entryPointParams_hmap_entries_2 : array<GridHashMapEntryGeneric_std430_1>;
@binding(12) @group(0) var<storage, read> entryPointParams_rigid_particles_pos_1 : array<Position_std430_0>;
@binding(13) @group(0) var<storage, read_write> entryPointParams_rigid_particle_needs_block_1 : array<atomic<u32>>;
@binding(14) @group(0) var<storage, read> entryPointParams_grid_3 : array<GridGeneric_std430_1>;
@binding(15) @group(0) var<storage, read> entryPointParams_hmap_entries_3 : array<GridHashMapEntryGeneric_std430_1>;
@binding(16) @group(0) var<storage, read> entryPointParams_particles_pos_1 : array<Position_std430_0>;
@binding(17) @group(0) var<uniform> entryPointParams_particles_len_1 : u32;
struct ActiveBlockHeaderGeneric_std430_1
{
@align(8) virtual_id_0 : BlockVirtualId_std430_0,
@align(8) first_particle_0 : u32,
@align(4) num_particles_0 : atomic<u32>,
};
@binding(18) @group(0) var<storage, read_write> entryPointParams_active_blocks_2 : array<ActiveBlockHeaderGeneric_std430_1>;
@binding(19) @group(0) var<storage, read> entryPointParams_grid_4 : array<GridGeneric_std430_1>;
@binding(20) @group(0) var<storage, read> entryPointParams_active_blocks_3 : array<ActiveBlockHeaderGeneric_std430_0>;
@binding(21) @group(0) var<storage, read_write> entryPointParams_scan_values_0 : array<u32>;
@binding(22) @group(0) var<storage, read> entryPointParams_grid_5 : array<GridGeneric_std430_1>;
@binding(23) @group(0) var<storage, read> entryPointParams_scan_values_1 : array<u32>;
@binding(24) @group(0) var<storage, read_write> entryPointParams_active_blocks_4 : array<ActiveBlockHeaderGeneric_std430_0>;
@binding(25) @group(0) var<storage, read> entryPointParams_grid_6 : array<GridGeneric_std430_1>;
@binding(26) @group(0) var<storage, read> entryPointParams_hmap_entries_4 : array<GridHashMapEntryGeneric_std430_1>;
@binding(27) @group(0) var<storage, read> entryPointParams_particles_pos_2 : array<Position_std430_0>;
@binding(28) @group(0) var<uniform> entryPointParams_particles_len_2 : u32;
@binding(29) @group(0) var<storage, read_write> entryPointParams_scan_values_2 : array<atomic<u32>>;
struct NodeLinkedListGeneric_std430_0
{
@align(4) head_0 : atomic<u32>,
@align(4) len_0 : atomic<u32>,
};
@binding(30) @group(0) var<storage, read_write> entryPointParams_nodes_linked_lists_0 : array<NodeLinkedListGeneric_std430_0>;
@binding(31) @group(0) var<storage, read_write> entryPointParams_particle_node_linked_lists_0 : array<u32>;
@binding(32) @group(0) var<storage, read_write> entryPointParams_sorted_particle_ids_0 : array<u32>;
@binding(33) @group(0) var<storage, read> entryPointParams_grid_7 : array<GridGeneric_std430_1>;
@binding(34) @group(0) var<storage, read> entryPointParams_hmap_entries_5 : array<GridHashMapEntryGeneric_std430_1>;
@binding(35) @group(0) var<storage, read> entryPointParams_rigid_particles_pos_2 : array<Position_std430_0>;
@binding(36) @group(0) var<storage, read_write> entryPointParams_rigid_nodes_linked_lists_0 : array<NodeLinkedListGeneric_std430_0>;
@binding(37) @group(0) var<storage, read_write> entryPointParams_rigid_particle_node_linked_lists_0 : array<u32>;
struct BlockVirtualId_0
{
id_0 : vec2<i32>,
};
fn BlockVirtualId_x24init_0( i_0 : vec2<i32>) -> BlockVirtualId_0
{
var _S1 : BlockVirtualId_0;
_S1.id_0 = i_0;
return _S1;
}
fn block_associated_to_point_0( cell_width_1 : f32, pt_1 : vec2<f32>) -> BlockVirtualId_0
{
var _S2 : vec2<f32> = floor((round(pt_1 / vec2<f32>(cell_width_1)) - vec2<f32>(1.0f)) / vec2<f32>(8.0f));
return BlockVirtualId_x24init_0(vec2<i32>(i32(_S2.x), i32(_S2.y)));
}
fn blocks_associated_to_block_0( block_0 : BlockVirtualId_0, _S3 : ptr<function, array<BlockVirtualId_0, i32(4)>>)
{
var _S4 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec2<i32>(i32(0), i32(1)));
var _S5 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec2<i32>(i32(1), i32(0)));
var _S6 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec2<i32>(i32(1), i32(1)));
(*_S3)[i32(0)] = BlockVirtualId_x24init_0(block_0.id_0);
(*_S3)[i32(1)] = _S4;
(*_S3)[i32(2)] = _S5;
(*_S3)[i32(3)] = _S6;
return;
}
fn blocks_associated_to_point_0( cell_width_2 : f32, pt_2 : vec2<f32>, _S7 : ptr<function, array<BlockVirtualId_0, i32(4)>>)
{
var _S8 : array<BlockVirtualId_0, i32(4)>;
blocks_associated_to_block_0(block_associated_to_point_0(cell_width_2, pt_2), &(_S8));
(*_S7) = _S8;
return;
}
fn pack_key_0( key_1 : BlockVirtualId_0) -> u32
{
return (((bitcast<u32>(key_1.id_0.x + i32(32767)) & (u32(65535)))) | (((((bitcast<u32>(key_1.id_0.y + i32(32767)) & (u32(65535)))) << (u32(16))))));
}
fn hash_0( packed_key_0 : u32) -> u32
{
var key_2 : u32 = packed_key_0 * u32(3432918353);
return ((((key_2 << (u32(15)))) | (((key_2 >> (u32(17))))))) * u32(461845907);
}
struct BlockHeaderId_0
{
id_1 : u32,
};
fn BlockHeaderId_x24init_0( i_1 : u32) -> BlockHeaderId_0
{
var _S9 : BlockHeaderId_0;
_S9.id_1 = i_1;
return _S9;
}
fn insertion_index_0( _S10 : u32, _S11 : BlockVirtualId_0) -> u32
{
var _S12 : u32 = pack_key_0(_S11);
var _S13 : u32 = _S10 - u32(1);
var slot_0 : u32 = ((hash_0(_S12)) & (_S13));
var k_0 : u32 = u32(0);
for(;;)
{
if(k_0 < _S10)
{
}
else
{
break;
}
var _S14 : u32 = atomicCompareExchangeWeak(&(entryPointParams_hmap_entries_0[slot_0].state_0), u32(4294967295), _S12).old_value;
if(_S14 == _S12)
{
return u32(4294967295);
}
else
{
if(_S14 == u32(4294967295))
{
entryPointParams_hmap_entries_0[slot_0].key_0.id_0 = _S11.id_0;
return slot_0;
}
}
var _S15 : u32 = (slot_0 + u32(1)) % _S10;
var _S16 : u32 = (_S15 & (_S13));
var _S17 : u32 = k_0 + u32(1);
slot_0 = _S16;
k_0 = _S17;
}
return u32(4294967295);
}
fn mark_block_as_active_0( _S18 : BlockVirtualId_0)
{
var _S19 : u32 = insertion_index_0(entryPointParams_grid_0[i32(0)].hmap_capacity_0, _S18);
if(_S19 != u32(4294967295))
{
var _S20 : u32 = atomicAdd(&(entryPointParams_grid_0[i32(0)].num_active_blocks_0), u32(1));
entryPointParams_active_blocks_0[_S20].virtual_id_0.id_0 = _S18.id_0;
entryPointParams_active_blocks_0[_S20].first_particle_0 = u32(0);
entryPointParams_active_blocks_0[_S20].num_particles_0 = u32(0);
entryPointParams_hmap_entries_0[_S19].value_0.id_1 = BlockHeaderId_x24init_0(_S20).id_1;
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id_0 : vec3<u32>)
{
var _S21 : u32 = invocation_id_0.x;
if(_S21 < entryPointParams_particles_len_0)
{
var _S22 : array<BlockVirtualId_0, i32(4)>;
blocks_associated_to_point_0(entryPointParams_grid_0[i32(0)].cell_width_0, entryPointParams_particles_pos_0[_S21].pt_0, &(_S22));
var _S23 : array<BlockVirtualId_0, i32(4)> = _S22;
var i_2 : u32 = u32(0);
for(;;)
{
if(i_2 < u32(4))
{
}
else
{
break;
}
mark_block_as_active_0(_S23[i_2]);
i_2 = i_2 + u32(1);
}
}
return;
}
fn getCount_0() -> i32
{
var _S24 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_0), 8);
return i32(_S24.x);
}
fn insertion_index_1( _S25 : u32, _S26 : BlockVirtualId_0) -> u32
{
var _S27 : u32 = pack_key_0(_S26);
var _S28 : u32 = _S25 - u32(1);
var slot_1 : u32 = ((hash_0(_S27)) & (_S28));
var k_1 : u32 = u32(0);
for(;;)
{
if(k_1 < _S25)
{
}
else
{
break;
}
var _S29 : u32 = atomicCompareExchangeWeak(&(entryPointParams_hmap_entries_1[slot_1].state_0), u32(4294967295), _S27).old_value;
if(_S29 == _S27)
{
return u32(4294967295);
}
else
{
if(_S29 == u32(4294967295))
{
entryPointParams_hmap_entries_1[slot_1].key_0.id_0 = _S26.id_0;
return slot_1;
}
}
var _S30 : u32 = (slot_1 + u32(1)) % _S25;
var _S31 : u32 = (_S30 & (_S28));
var _S32 : u32 = k_1 + u32(1);
slot_1 = _S31;
k_1 = _S32;
}
return u32(4294967295);
}
fn mark_block_as_active_1( _S33 : BlockVirtualId_0)
{
var _S34 : u32 = insertion_index_1(entryPointParams_grid_1[i32(0)].hmap_capacity_0, _S33);
if(_S34 != u32(4294967295))
{
var _S35 : u32 = atomicAdd(&(entryPointParams_grid_1[i32(0)].num_active_blocks_0), u32(1));
entryPointParams_active_blocks_1[_S35].virtual_id_0.id_0 = _S33.id_0;
entryPointParams_active_blocks_1[_S35].first_particle_0 = u32(0);
entryPointParams_active_blocks_1[_S35].num_particles_0 = u32(0);
entryPointParams_hmap_entries_1[_S34].value_0.id_1 = BlockHeaderId_x24init_0(_S35).id_1;
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn touch_rigid_particle_blocks(@builtin(global_invocation_id) invocation_id_1 : vec3<u32>)
{
var _S36 : u32 = invocation_id_1.x;
var _S37 : i32 = getCount_0();
if(_S36 < u32(_S37))
{
var _S38 : f32 = entryPointParams_grid_1[i32(0)].cell_width_0;
if(((entryPointParams_rigid_particle_needs_block_0[_S36 / u32(32)] & (((u32(1) << ((_S36 % u32(32)))))))) != u32(0))
{
mark_block_as_active_1(block_associated_to_point_0(_S38, entryPointParams_rigid_particles_pos_0[_S36].pt_0));
}
}
return;
}
fn getCount_1() -> i32
{
var _S39 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_1), 8);
return i32(_S39.x);
}
fn find_block_header_id_0( _S40 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S41 : u32 = pack_key_0(_S40);
var slot_2 : u32 = ((hash_0(_S41)) & ((entryPointParams_grid_2[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S42 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_2[slot_2];
if((_S42.state_0) == _S41)
{
var _S43 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_2[slot_2].value_0.id_1 );
return _S43;
}
else
{
if((_S42.state_0) == u32(4294967295))
{
break;
}
}
slot_2 = ((slot_2 + u32(1)) & ((entryPointParams_grid_2[i32(0)].hmap_capacity_0 - u32(1))));
}
return BlockHeaderId_x24init_0(u32(4294967295));
}
@compute
@workgroup_size(64, 1, 1)
fn mark_rigid_particles_needing_block(@builtin(global_invocation_id) invocation_id_2 : vec3<u32>)
{
var _S44 : bool;
var _S45 : u32 = invocation_id_2.x;
var _S46 : i32 = getCount_1();
if(_S45 < u32(_S46))
{
var _S47 : array<BlockVirtualId_0, i32(4)>;
blocks_associated_to_point_0(entryPointParams_grid_2[i32(0)].cell_width_0, entryPointParams_rigid_particles_pos_1[_S45].pt_0, &(_S47));
var _S48 : array<BlockVirtualId_0, i32(4)> = _S47;
var i_3 : u32 = u32(0);
for(;;)
{
var _S49 : bool = i_3 < u32(4);
_S44 = _S49;
if(_S49)
{
}
else
{
break;
}
if((find_block_header_id_0(_S48[i_3]).id_1) != u32(4294967295))
{
break;
}
i_3 = i_3 + u32(1);
}
var _S50 : u32 = _S45 / u32(32);
var _S51 : u32 = (u32(1) << ((_S45 % u32(32))));
var _S52 : bool;
if(i_3 > u32(0))
{
_S52 = _S44;
}
else
{
_S52 = false;
}
if(_S52)
{
var _S53 : u32 = atomicOr(&(entryPointParams_rigid_particle_needs_block_1[_S50]), _S51);
}
else
{
var _S54 : u32 = atomicAnd(&(entryPointParams_rigid_particle_needs_block_1[_S50]), ~_S51);
}
}
return;
}
fn find_block_header_id_1( _S55 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S56 : u32 = pack_key_0(_S55);
var slot_3 : u32 = ((hash_0(_S56)) & ((entryPointParams_grid_3[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S57 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_3[slot_3];
if((_S57.state_0) == _S56)
{
var _S58 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_3[slot_3].value_0.id_1 );
return _S58;
}
else
{
if((_S57.state_0) == u32(4294967295))
{
break;
}
}
slot_3 = ((slot_3 + u32(1)) & ((entryPointParams_grid_3[i32(0)].hmap_capacity_0 - u32(1))));
}
return BlockHeaderId_x24init_0(u32(4294967295));
}
@compute
@workgroup_size(64, 1, 1)
fn update_block_particle_count(@builtin(global_invocation_id) invocation_id_3 : vec3<u32>)
{
var _S59 : u32 = invocation_id_3.x;
if(_S59 < entryPointParams_particles_len_1)
{
var _S60 : u32 = atomicAdd(&(entryPointParams_active_blocks_2[find_block_header_id_1(block_associated_to_point_0(entryPointParams_grid_3[i32(0)].cell_width_0, entryPointParams_particles_pos_1[_S59].pt_0)).id_1].num_particles_0), u32(1));
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn copy_particles_len_to_scan_value(@builtin(global_invocation_id) invocation_id_4 : vec3<u32>)
{
var _S61 : u32 = invocation_id_4.x;
if(_S61 < (entryPointParams_grid_4[i32(0)].num_active_blocks_0))
{
entryPointParams_scan_values_0[_S61] = entryPointParams_active_blocks_3[_S61].num_particles_0;
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn copy_scan_values_to_first_particles(@builtin(global_invocation_id) invocation_id_5 : vec3<u32>)
{
var _S62 : u32 = invocation_id_5.x;
if(_S62 < (entryPointParams_grid_5[i32(0)].num_active_blocks_0))
{
entryPointParams_active_blocks_4[_S62].first_particle_0 = entryPointParams_scan_values_1[_S62];
}
return;
}
fn associated_cell_index_in_block_off_by_one_0( part_pos_0 : ptr<function, Position_std430_0>, cell_width_3 : f32) -> vec2<u32>
{
var _S63 : vec2<f32> = round((*part_pos_0).pt_0 / vec2<f32>(cell_width_3)) - vec2<f32>(1.0f);
var _S64 : vec2<f32> = vec2<f32>(8.0f);
return vec2<u32>(_S63 - floor(_S63 / _S64) * _S64);
}
struct BlockPhysicalId_0
{
id_2 : u32,
};
fn BlockPhysicalId_x24init_0( i_4 : u32) -> BlockPhysicalId_0
{
var _S65 : BlockPhysicalId_0;
_S65.id_2 = i_4;
return _S65;
}
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_5 : u32) -> NodePhysicalId_0
{
var _S66 : NodePhysicalId_0;
_S66.id_3 = i_5;
return _S66;
}
fn node_id_0( pid_0 : BlockPhysicalId_0, shift_in_block_0 : vec2<u32>) -> NodePhysicalId_0
{
return NodePhysicalId_x24init_0(pid_0.id_2 + shift_in_block_0.x + shift_in_block_0.y * u32(8));
}
fn find_block_header_id_2( _S67 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S68 : u32 = pack_key_0(_S67);
var slot_4 : u32 = ((hash_0(_S68)) & ((entryPointParams_grid_6[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S69 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_4[slot_4];
if((_S69.state_0) == _S68)
{
var _S70 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_4[slot_4].value_0.id_1 );
return _S70;
}
else
{
if((_S69.state_0) == u32(4294967295))
{
break;
}
}
slot_4 = ((slot_4 + u32(1)) & ((entryPointParams_grid_6[i32(0)].hmap_capacity_0 - u32(1))));
}
return BlockHeaderId_x24init_0(u32(4294967295));
}
@compute
@workgroup_size(64, 1, 1)
fn finalize_particles_sort(@builtin(global_invocation_id) invocation_id_6 : vec3<u32>)
{
var _S71 : u32 = invocation_id_6.x;
if(_S71 < entryPointParams_particles_len_2)
{
var _S72 : GridGeneric_std430_1 = entryPointParams_grid_6[i32(0)];
var _S73 : Position_std430_0 = entryPointParams_particles_pos_2[_S71];
var _S74 : BlockHeaderId_0 = find_block_header_id_2(block_associated_to_point_0(_S72.cell_width_0, _S73.pt_0));
var _S75 : u32 = atomicAdd(&(entryPointParams_scan_values_2[_S74.id_1]), u32(1));
entryPointParams_sorted_particle_ids_0[_S75] = _S71;
var _S76 : vec2<u32> = associated_cell_index_in_block_off_by_one_0(&(_S73), _S72.cell_width_0);
var _S77 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S74), _S76);
var _S78 : u32 = atomicExchange(&(entryPointParams_nodes_linked_lists_0[_S77.id_3].head_0), _S71);
var _S79 : u32 = atomicAdd(&(entryPointParams_nodes_linked_lists_0[_S77.id_3].len_0), u32(1));
entryPointParams_particle_node_linked_lists_0[_S71] = _S78;
}
return;
}
fn getCount_2() -> i32
{
var _S80 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_2), 8);
return i32(_S80.x);
}
fn find_block_header_id_3( _S81 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S82 : u32 = pack_key_0(_S81);
var slot_5 : u32 = ((hash_0(_S82)) & ((entryPointParams_grid_7[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S83 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_5[slot_5];
if((_S83.state_0) == _S82)
{
var _S84 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_5[slot_5].value_0.id_1 );
return _S84;
}
else
{
if((_S83.state_0) == u32(4294967295))
{
break;
}
}
slot_5 = ((slot_5 + u32(1)) & ((entryPointParams_grid_7[i32(0)].hmap_capacity_0 - u32(1))));
}
return BlockHeaderId_x24init_0(u32(4294967295));
}
@compute
@workgroup_size(64, 1, 1)
fn sort_rigid_particles(@builtin(global_invocation_id) invocation_id_7 : vec3<u32>)
{
var _S85 : u32 = invocation_id_7.x;
var _S86 : i32 = getCount_2();
if(_S85 < u32(_S86))
{
var _S87 : GridGeneric_std430_1 = entryPointParams_grid_7[i32(0)];
var _S88 : Position_std430_0 = entryPointParams_rigid_particles_pos_2[_S85];
var _S89 : BlockHeaderId_0 = find_block_header_id_3(block_associated_to_point_0(_S87.cell_width_0, _S88.pt_0));
if((_S89.id_1) != u32(4294967295))
{
var _S90 : vec2<u32> = associated_cell_index_in_block_off_by_one_0(&(_S88), _S87.cell_width_0);
var _S91 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S89), _S90);
var _S92 : u32 = atomicExchange(&(entryPointParams_rigid_nodes_linked_lists_0[_S91.id_3].head_0), _S85);
var _S93 : u32 = atomicAdd(&(entryPointParams_rigid_nodes_linked_lists_0[_S91.id_3].len_0), u32(1));
entryPointParams_rigid_particle_node_linked_lists_0[_S85] = _S92;
}
}
return;
}