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