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(16) id_0 : vec3<i32>,
};
struct BlockHeaderId_std430_0
{
@align(4) id_1 : u32,
};
struct GridHashMapEntryGeneric_std430_0
{
@align(16) state_0 : atomic<u32>,
@align(16) key_0 : BlockVirtualId_std430_0,
@align(16) 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(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_write> entryPointParams_active_blocks_0 : array<ActiveBlockHeaderGeneric_std430_0>;
struct Position_std430_0
{
@align(16) pt_0 : vec3<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(16) state_0 : u32,
@align(16) key_0 : BlockVirtualId_std430_0,
@align(16) 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(16) virtual_id_0 : BlockVirtualId_std430_0,
@align(16) 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 : vec3<i32>,
};
fn BlockVirtualId_x24init_0( i_0 : vec3<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 : vec3<f32>) -> BlockVirtualId_0
{
var _S2 : vec3<f32> = floor((round(pt_1 / vec3<f32>(cell_width_1)) - vec3<f32>(1.0f)) / vec3<f32>(4.0f));
return BlockVirtualId_x24init_0(vec3<i32>(i32(_S2.x), i32(_S2.y), i32(_S2.z)));
}
fn blocks_associated_to_block_0( block_0 : BlockVirtualId_0, _S3 : ptr<function, array<BlockVirtualId_0, i32(8)>>)
{
var _S4 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(0), i32(0), i32(1)));
var _S5 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(0), i32(1), i32(0)));
var _S6 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(0), i32(1), i32(1)));
var _S7 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(1), i32(0), i32(0)));
var _S8 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(1), i32(0), i32(1)));
var _S9 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(1), i32(1), i32(0)));
var _S10 : BlockVirtualId_0 = BlockVirtualId_x24init_0(block_0.id_0 + vec3<i32>(i32(1), 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;
(*_S3)[i32(4)] = _S7;
(*_S3)[i32(5)] = _S8;
(*_S3)[i32(6)] = _S9;
(*_S3)[i32(7)] = _S10;
return;
}
fn blocks_associated_to_point_0( cell_width_2 : f32, pt_2 : vec3<f32>, _S11 : ptr<function, array<BlockVirtualId_0, i32(8)>>)
{
var _S12 : array<BlockVirtualId_0, i32(8)>;
blocks_associated_to_block_0(block_associated_to_point_0(cell_width_2, pt_2), &(_S12));
(*_S11) = _S12;
return;
}
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 _S13 : BlockHeaderId_0;
_S13.id_1 = i_1;
return _S13;
}
fn insertion_index_0( _S14 : u32, _S15 : BlockVirtualId_0) -> u32
{
var _S16 : u32 = pack_key_0(_S15);
var _S17 : u32 = _S14 - u32(1);
var slot_0 : u32 = ((hash_0(_S16)) & (_S17));
var k_0 : u32 = u32(0);
for(;;)
{
if(k_0 < _S14)
{
}
else
{
break;
}
var _S18 : u32 = atomicCompareExchangeWeak(&(entryPointParams_hmap_entries_0[slot_0].state_0), u32(4294967295), _S16).old_value;
if(_S18 == _S16)
{
return u32(4294967295);
}
else
{
if(_S18 == u32(4294967295))
{
entryPointParams_hmap_entries_0[slot_0].key_0.id_0 = _S15.id_0;
return slot_0;
}
}
var _S19 : u32 = (slot_0 + u32(1)) % _S14;
var _S20 : u32 = (_S19 & (_S17));
var _S21 : u32 = k_0 + u32(1);
slot_0 = _S20;
k_0 = _S21;
}
return u32(4294967295);
}
fn mark_block_as_active_0( _S22 : BlockVirtualId_0)
{
var _S23 : u32 = insertion_index_0(entryPointParams_grid_0[i32(0)].hmap_capacity_0, _S22);
if(_S23 != u32(4294967295))
{
var _S24 : u32 = atomicAdd(&(entryPointParams_grid_0[i32(0)].num_active_blocks_0), u32(1));
entryPointParams_active_blocks_0[_S24].virtual_id_0.id_0 = _S22.id_0;
entryPointParams_active_blocks_0[_S24].first_particle_0 = u32(0);
entryPointParams_active_blocks_0[_S24].num_particles_0 = u32(0);
entryPointParams_hmap_entries_0[_S23].value_0.id_1 = BlockHeaderId_x24init_0(_S24).id_1;
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id_0 : vec3<u32>)
{
var _S25 : u32 = invocation_id_0.x;
if(_S25 < entryPointParams_particles_len_0)
{
var _S26 : array<BlockVirtualId_0, i32(8)>;
blocks_associated_to_point_0(entryPointParams_grid_0[i32(0)].cell_width_0, entryPointParams_particles_pos_0[_S25].pt_0, &(_S26));
var _S27 : array<BlockVirtualId_0, i32(8)> = _S26;
var i_2 : u32 = u32(0);
for(;;)
{
if(i_2 < u32(8))
{
}
else
{
break;
}
mark_block_as_active_0(_S27[i_2]);
i_2 = i_2 + u32(1);
}
}
return;
}
fn getCount_0() -> i32
{
var _S28 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_0), 16);
return i32(_S28.x);
}
fn insertion_index_1( _S29 : u32, _S30 : BlockVirtualId_0) -> u32
{
var _S31 : u32 = pack_key_0(_S30);
var _S32 : u32 = _S29 - u32(1);
var slot_1 : u32 = ((hash_0(_S31)) & (_S32));
var k_1 : u32 = u32(0);
for(;;)
{
if(k_1 < _S29)
{
}
else
{
break;
}
var _S33 : u32 = atomicCompareExchangeWeak(&(entryPointParams_hmap_entries_1[slot_1].state_0), u32(4294967295), _S31).old_value;
if(_S33 == _S31)
{
return u32(4294967295);
}
else
{
if(_S33 == u32(4294967295))
{
entryPointParams_hmap_entries_1[slot_1].key_0.id_0 = _S30.id_0;
return slot_1;
}
}
var _S34 : u32 = (slot_1 + u32(1)) % _S29;
var _S35 : u32 = (_S34 & (_S32));
var _S36 : u32 = k_1 + u32(1);
slot_1 = _S35;
k_1 = _S36;
}
return u32(4294967295);
}
fn mark_block_as_active_1( _S37 : BlockVirtualId_0)
{
var _S38 : u32 = insertion_index_1(entryPointParams_grid_1[i32(0)].hmap_capacity_0, _S37);
if(_S38 != u32(4294967295))
{
var _S39 : u32 = atomicAdd(&(entryPointParams_grid_1[i32(0)].num_active_blocks_0), u32(1));
entryPointParams_active_blocks_1[_S39].virtual_id_0.id_0 = _S37.id_0;
entryPointParams_active_blocks_1[_S39].first_particle_0 = u32(0);
entryPointParams_active_blocks_1[_S39].num_particles_0 = u32(0);
entryPointParams_hmap_entries_1[_S38].value_0.id_1 = BlockHeaderId_x24init_0(_S39).id_1;
}
return;
}
@compute
@workgroup_size(64, 1, 1)
fn touch_rigid_particle_blocks(@builtin(global_invocation_id) invocation_id_1 : vec3<u32>)
{
var _S40 : u32 = invocation_id_1.x;
var _S41 : i32 = getCount_0();
if(_S40 < u32(_S41))
{
var _S42 : f32 = entryPointParams_grid_1[i32(0)].cell_width_0;
if(((entryPointParams_rigid_particle_needs_block_0[_S40 / u32(32)] & (((u32(1) << ((_S40 % u32(32)))))))) != u32(0))
{
mark_block_as_active_1(block_associated_to_point_0(_S42, entryPointParams_rigid_particles_pos_0[_S40].pt_0));
}
}
return;
}
fn getCount_1() -> i32
{
var _S43 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_1), 16);
return i32(_S43.x);
}
fn find_block_header_id_0( _S44 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S45 : u32 = pack_key_0(_S44);
var slot_2 : u32 = ((hash_0(_S45)) & ((entryPointParams_grid_2[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S46 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_2[slot_2];
if((_S46.state_0) == _S45)
{
var _S47 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_2[slot_2].value_0.id_1 );
return _S47;
}
else
{
if((_S46.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 _S48 : bool;
var _S49 : u32 = invocation_id_2.x;
var _S50 : i32 = getCount_1();
if(_S49 < u32(_S50))
{
var _S51 : array<BlockVirtualId_0, i32(8)>;
blocks_associated_to_point_0(entryPointParams_grid_2[i32(0)].cell_width_0, entryPointParams_rigid_particles_pos_1[_S49].pt_0, &(_S51));
var _S52 : array<BlockVirtualId_0, i32(8)> = _S51;
var i_3 : u32 = u32(0);
for(;;)
{
var _S53 : bool = i_3 < u32(8);
_S48 = _S53;
if(_S53)
{
}
else
{
break;
}
if((find_block_header_id_0(_S52[i_3]).id_1) != u32(4294967295))
{
break;
}
i_3 = i_3 + u32(1);
}
var _S54 : u32 = _S49 / u32(32);
var _S55 : u32 = (u32(1) << ((_S49 % u32(32))));
var _S56 : bool;
if(i_3 > u32(0))
{
_S56 = _S48;
}
else
{
_S56 = false;
}
if(_S56)
{
var _S57 : u32 = atomicOr(&(entryPointParams_rigid_particle_needs_block_1[_S54]), _S55);
}
else
{
var _S58 : u32 = atomicAnd(&(entryPointParams_rigid_particle_needs_block_1[_S54]), ~_S55);
}
}
return;
}
fn find_block_header_id_1( _S59 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S60 : u32 = pack_key_0(_S59);
var slot_3 : u32 = ((hash_0(_S60)) & ((entryPointParams_grid_3[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S61 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_3[slot_3];
if((_S61.state_0) == _S60)
{
var _S62 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_3[slot_3].value_0.id_1 );
return _S62;
}
else
{
if((_S61.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 _S63 : u32 = invocation_id_3.x;
if(_S63 < entryPointParams_particles_len_1)
{
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[_S63].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;
}
fn associated_cell_index_in_block_off_by_one_0( part_pos_0 : ptr<function, Position_std430_0>, cell_width_3 : f32) -> vec3<u32>
{
var _S67 : vec3<f32> = round((*part_pos_0).pt_0 / vec3<f32>(cell_width_3)) - vec3<f32>(1.0f);
var _S68 : vec3<f32> = vec3<f32>(4.0f);
return vec3<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 : 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));
}
fn find_block_header_id_2( _S71 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S72 : u32 = pack_key_0(_S71);
var slot_4 : u32 = ((hash_0(_S72)) & ((entryPointParams_grid_6[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S73 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_4[slot_4];
if((_S73.state_0) == _S72)
{
var _S74 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_4[slot_4].value_0.id_1 );
return _S74;
}
else
{
if((_S73.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 _S75 : u32 = invocation_id_6.x;
if(_S75 < entryPointParams_particles_len_2)
{
var _S76 : GridGeneric_std430_1 = entryPointParams_grid_6[i32(0)];
var _S77 : Position_std430_0 = entryPointParams_particles_pos_2[_S75];
var _S78 : BlockHeaderId_0 = find_block_header_id_2(block_associated_to_point_0(_S76.cell_width_0, _S77.pt_0));
var _S79 : u32 = atomicAdd(&(entryPointParams_scan_values_2[_S78.id_1]), u32(1));
entryPointParams_sorted_particle_ids_0[_S79] = _S75;
var _S80 : vec3<u32> = associated_cell_index_in_block_off_by_one_0(&(_S77), _S76.cell_width_0);
var _S81 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S78), _S80);
var _S82 : u32 = atomicExchange(&(entryPointParams_nodes_linked_lists_0[_S81.id_3].head_0), _S75);
var _S83 : u32 = atomicAdd(&(entryPointParams_nodes_linked_lists_0[_S81.id_3].len_0), u32(1));
entryPointParams_particle_node_linked_lists_0[_S75] = _S82;
}
return;
}
fn getCount_2() -> i32
{
var _S84 : vec2<u32> = vec2<u32>(arrayLength(&entryPointParams_rigid_particles_pos_2), 16);
return i32(_S84.x);
}
fn find_block_header_id_3( _S85 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S86 : u32 = pack_key_0(_S85);
var slot_5 : u32 = ((hash_0(_S86)) & ((entryPointParams_grid_7[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S87 : GridHashMapEntryGeneric_std430_1 = entryPointParams_hmap_entries_5[slot_5];
if((_S87.state_0) == _S86)
{
var _S88 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_5[slot_5].value_0.id_1 );
return _S88;
}
else
{
if((_S87.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 _S89 : u32 = invocation_id_7.x;
var _S90 : i32 = getCount_2();
if(_S89 < u32(_S90))
{
var _S91 : GridGeneric_std430_1 = entryPointParams_grid_7[i32(0)];
var _S92 : Position_std430_0 = entryPointParams_rigid_particles_pos_2[_S89];
var _S93 : BlockHeaderId_0 = find_block_header_id_3(block_associated_to_point_0(_S91.cell_width_0, _S92.pt_0));
if((_S93.id_1) != u32(4294967295))
{
var _S94 : vec3<u32> = associated_cell_index_in_block_off_by_one_0(&(_S92), _S91.cell_width_0);
var _S95 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S93), _S94);
var _S96 : u32 = atomicExchange(&(entryPointParams_rigid_nodes_linked_lists_0[_S95.id_3].head_0), _S89);
var _S97 : u32 = atomicAdd(&(entryPointParams_rigid_nodes_linked_lists_0[_S95.id_3].len_0), u32(1));
entryPointParams_rigid_particle_node_linked_lists_0[_S89] = _S96;
}
}
return;
}