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