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(8) id_0 : vec2<i32>,
};
struct BlockHeaderId_std430_0
{
@align(4) id_1 : u32,
};
struct GridHashMapEntryGeneric_std430_0
{
@align(8) state_0 : u32,
@align(8) key_0 : BlockVirtualId_std430_0,
@align(8) 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(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> 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_rigid_nodes_linked_lists_0 : array<NodeLinkedListGeneric_std430_0>;
@binding(4) @group(0) var<storage, read> entryPointParams_particle_node_linked_lists_0 : array<u32>;
@binding(5) @group(0) var<storage, read> entryPointParams_collider_vertices_0 : array<vec2<f32>>;
struct RigidParticleIndices_std430_0
{
@align(8) segment_0 : vec2<u32>,
@align(8) collider_0 : u32,
};
@binding(6) @group(0) var<storage, read> entryPointParams_rigid_particle_indices_0 : array<RigidParticleIndices_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 : vec3<f32>,
@align(4) cdf_0 : NodeCdf_std430_0,
};
@binding(7) @group(0) var<storage, read_write> entryPointParams_nodes_0 : array<Node_std430_0>;
var<workgroup> max_linked_list_length_0 : atomic<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 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 _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 : vec2<u32>) -> NodePhysicalId_0
{
return NodePhysicalId_x24init_0(pid_0.id_2 + shift_in_block_0.x + shift_in_block_0.y * u32(8));
}
var<workgroup> max_linked_list_length_uniform_0 : u32;
fn flatten_shared_index_0( x_0 : u32, y_0 : u32) -> u32
{
return x_0 - u32(6) + (y_0 - u32(6)) * u32(10);
}
struct SharedNode_0
{
particle_id_0 : u32,
global_id_0 : u32,
};
var<workgroup> shared_nodes_0 : array<SharedNode_0, i32(100)>;
var<workgroup> shared_collider_ids_0 : array<u32, i32(100)>;
struct Segment_0
{
a_0 : vec2<f32>,
b_0 : vec2<f32>,
};
var<workgroup> shared_primitives_0 : array<Segment_0, i32(100)>;
fn Segment_x24init_0( a_1 : vec2<f32>, b_1 : vec2<f32>) -> Segment_0
{
var _S5 : Segment_0;
_S5.a_0 = a_1;
_S5.b_0 = b_1;
return _S5;
}
fn flatten_shared_shift_0( x_1 : u32, y_1 : u32) -> u32
{
return x_1 + y_1 * u32(10);
}
struct NodeCdf_0
{
distance_0 : f32,
affinities_0 : u32,
closest_id_0 : u32,
};
fn NodeCdf_x24init_0( distance_1 : f32, affinities_1 : u32, closest_id_1 : u32) -> NodeCdf_0
{
var _S6 : NodeCdf_0;
_S6.distance_0 = distance_1;
_S6.affinities_0 = affinities_1;
_S6.closest_id_0 = closest_id_1;
return _S6;
}
fn Segment_project_local_point_0( this_0 : Segment_0, pt_0 : vec2<f32>) -> vec2<f32>
{
var _S7 : vec2<f32> = this_0.b_0 - this_0.a_0;
var _S8 : f32 = dot(_S7, pt_0 - this_0.a_0);
var _S9 : f32 = dot(_S7, _S7);
if(_S8 <= 0.0f)
{
return this_0.a_0;
}
else
{
if(_S8 >= _S9)
{
return this_0.b_0;
}
else
{
return this_0.a_0 + _S7 * vec2<f32>((_S8 / _S9));
}
}
}
fn p2g_step_0( packed_cell_index_in_block_0 : u32, cell_width_1 : f32, cell_pos_0 : vec2<f32>) -> NodeCdf_0
{
var _S10 : u32;
var _S11 : u32 = flatten_shared_shift_0(u32(2), u32(2));
var result_0 : NodeCdf_0 = NodeCdf_x24init_0(1.0e+10f, u32(0), u32(4294967295));
for(;;)
{
var _S12 : bool;
for(;;)
{
for(;;)
{
var _S13 : u32 = packed_cell_index_in_block_0 - _S11;
_S10 = _S13;
var _S14 : u32 = _S13 + u32(22);
var _S15 : u32 = shared_collider_ids_0[_S14];
if((shared_collider_ids_0[_S14]) == u32(4294967295))
{
break;
}
var _S16 : Segment_0 = shared_primitives_0[_S14];
var _S17 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S14], cell_pos_0);
if((any((_S17 != (shared_primitives_0[_S14].a_0)))))
{
_S12 = (any((_S17 != (_S16.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S18 : vec2<f32> = cell_pos_0 - _S17;
var _S19 : f32 = length(_S18);
var _S20 : vec2<f32> = _S16.b_0 - _S16.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S15))) | (((u32((dot(_S18, vec2<f32>(- _S20.y, _S20.x))) < 0.0f) << ((_S15 + u32(16))))))))));
if(_S19 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S19);
result_0.closest_id_0 = _S15;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S21 : u32 = _S10 + u32(2);
var _S22 : u32 = shared_collider_ids_0[_S21];
if((shared_collider_ids_0[_S21]) == u32(4294967295))
{
break;
}
var _S23 : Segment_0 = shared_primitives_0[_S21];
var _S24 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S21], cell_pos_0);
if((any((_S24 != (shared_primitives_0[_S21].a_0)))))
{
_S12 = (any((_S24 != (_S23.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S25 : vec2<f32> = cell_pos_0 - _S24;
var _S26 : f32 = length(_S25);
var _S27 : vec2<f32> = _S23.b_0 - _S23.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S22))) | (((u32((dot(_S25, vec2<f32>(- _S27.y, _S27.x))) < 0.0f) << ((_S22 + u32(16))))))))));
if(_S26 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S26);
result_0.closest_id_0 = _S22;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S28 : u32 = _S10 + u32(12);
var _S29 : u32 = shared_collider_ids_0[_S28];
if((shared_collider_ids_0[_S28]) == u32(4294967295))
{
break;
}
var _S30 : Segment_0 = shared_primitives_0[_S28];
var _S31 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S28], cell_pos_0);
if((any((_S31 != (shared_primitives_0[_S28].a_0)))))
{
_S12 = (any((_S31 != (_S30.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S32 : vec2<f32> = cell_pos_0 - _S31;
var _S33 : f32 = length(_S32);
var _S34 : vec2<f32> = _S30.b_0 - _S30.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S29))) | (((u32((dot(_S32, vec2<f32>(- _S34.y, _S34.x))) < 0.0f) << ((_S29 + u32(16))))))))));
if(_S33 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S33);
result_0.closest_id_0 = _S29;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S35 : u32 = _S10 + u32(20);
var _S36 : u32 = shared_collider_ids_0[_S35];
if((shared_collider_ids_0[_S35]) == u32(4294967295))
{
break;
}
var _S37 : Segment_0 = shared_primitives_0[_S35];
var _S38 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S35], cell_pos_0);
if((any((_S38 != (shared_primitives_0[_S35].a_0)))))
{
_S12 = (any((_S38 != (_S37.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S39 : vec2<f32> = cell_pos_0 - _S38;
var _S40 : f32 = length(_S39);
var _S41 : vec2<f32> = _S37.b_0 - _S37.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S36))) | (((u32((dot(_S39, vec2<f32>(- _S41.y, _S41.x))) < 0.0f) << ((_S36 + u32(16))))))))));
if(_S40 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S40);
result_0.closest_id_0 = _S36;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S42 : u32 = shared_collider_ids_0[_S10];
if((shared_collider_ids_0[_S10]) == u32(4294967295))
{
break;
}
var _S43 : Segment_0 = shared_primitives_0[_S10];
var _S44 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S10], cell_pos_0);
if((any((_S44 != (shared_primitives_0[_S10].a_0)))))
{
_S12 = (any((_S44 != (_S43.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S45 : vec2<f32> = cell_pos_0 - _S44;
var _S46 : f32 = length(_S45);
var _S47 : vec2<f32> = _S43.b_0 - _S43.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S42))) | (((u32((dot(_S45, vec2<f32>(- _S47.y, _S47.x))) < 0.0f) << ((_S42 + u32(16))))))))));
if(_S46 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S46);
result_0.closest_id_0 = _S42;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S48 : u32 = _S10 + u32(10);
var _S49 : u32 = shared_collider_ids_0[_S48];
if((shared_collider_ids_0[_S48]) == u32(4294967295))
{
break;
}
var _S50 : Segment_0 = shared_primitives_0[_S48];
var _S51 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S48], cell_pos_0);
if((any((_S51 != (shared_primitives_0[_S48].a_0)))))
{
_S12 = (any((_S51 != (_S50.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S52 : vec2<f32> = cell_pos_0 - _S51;
var _S53 : f32 = length(_S52);
var _S54 : vec2<f32> = _S50.b_0 - _S50.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S49))) | (((u32((dot(_S52, vec2<f32>(- _S54.y, _S54.x))) < 0.0f) << ((_S49 + u32(16))))))))));
if(_S53 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S53);
result_0.closest_id_0 = _S49;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S55 : u32 = _S10 + u32(21);
var _S56 : u32 = shared_collider_ids_0[_S55];
if((shared_collider_ids_0[_S55]) == u32(4294967295))
{
break;
}
var _S57 : Segment_0 = shared_primitives_0[_S55];
var _S58 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S55], cell_pos_0);
if((any((_S58 != (shared_primitives_0[_S55].a_0)))))
{
_S12 = (any((_S58 != (_S57.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S59 : vec2<f32> = cell_pos_0 - _S58;
var _S60 : f32 = length(_S59);
var _S61 : vec2<f32> = _S57.b_0 - _S57.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S56))) | (((u32((dot(_S59, vec2<f32>(- _S61.y, _S61.x))) < 0.0f) << ((_S56 + u32(16))))))))));
if(_S60 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S60);
result_0.closest_id_0 = _S56;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S62 : u32 = _S10 + u32(1);
var _S63 : u32 = shared_collider_ids_0[_S62];
if((shared_collider_ids_0[_S62]) == u32(4294967295))
{
break;
}
var _S64 : Segment_0 = shared_primitives_0[_S62];
var _S65 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S62], cell_pos_0);
if((any((_S65 != (shared_primitives_0[_S62].a_0)))))
{
_S12 = (any((_S65 != (_S64.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S66 : vec2<f32> = cell_pos_0 - _S65;
var _S67 : f32 = length(_S66);
var _S68 : vec2<f32> = _S64.b_0 - _S64.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S63))) | (((u32((dot(_S66, vec2<f32>(- _S68.y, _S68.x))) < 0.0f) << ((_S63 + u32(16))))))))));
if(_S67 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S67);
result_0.closest_id_0 = _S63;
}
}
break;
}
break;
}
for(;;)
{
for(;;)
{
var _S69 : u32 = _S10 + u32(11);
var _S70 : u32 = shared_collider_ids_0[_S69];
if((shared_collider_ids_0[_S69]) == u32(4294967295))
{
break;
}
var _S71 : Segment_0 = shared_primitives_0[_S69];
var _S72 : vec2<f32> = Segment_project_local_point_0(shared_primitives_0[_S69], cell_pos_0);
if((any((_S72 != (shared_primitives_0[_S69].a_0)))))
{
_S12 = (any((_S72 != (_S71.b_0))));
}
else
{
_S12 = false;
}
if(_S12)
{
var _S73 : vec2<f32> = cell_pos_0 - _S72;
var _S74 : f32 = length(_S73);
var _S75 : vec2<f32> = _S71.b_0 - _S71.a_0;
result_0.affinities_0 = ((result_0.affinities_0) | (((((u32(1) << (_S70))) | (((u32((dot(_S73, vec2<f32>(- _S75.y, _S75.x))) < 0.0f) << ((_S70 + u32(16))))))))));
if(_S74 < (result_0.distance_0))
{
result_0.distance_0 = min(result_0.distance_0, _S74);
result_0.closest_id_0 = _S70;
}
}
break;
}
break;
}
break;
}
return result_0;
}
fn find_block_header_id_0( _S76 : BlockVirtualId_0) -> BlockHeaderId_0
{
var _S77 : u32 = pack_key_0(_S76);
var slot_0 : u32 = ((hash_0(_S77)) & ((entryPointParams_grid_0[i32(0)].hmap_capacity_0 - u32(1))));
for(;;)
{
var _S78 : GridHashMapEntryGeneric_std430_0 = entryPointParams_hmap_entries_0[slot_0];
if((_S78.state_0) == _S77)
{
var _S79 : BlockHeaderId_0 = BlockHeaderId_0( entryPointParams_hmap_entries_0[slot_0].value_0.id_1 );
return _S79;
}
else
{
if((_S78.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( _S80 : vec3<u32>, _S81 : ptr<function, BlockVirtualId_std430_0>, _S82 : u32)
{
var _S83 : vec2<i32> = (*_S81).id_0 - vec2<i32>(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(0))
{
}
else
{
break;
}
var _S84 : bool;
if(i_4 == u32(0))
{
_S84 = (_S80.x) < u32(6);
}
else
{
_S84 = false;
}
var _S85 : bool;
if(_S84)
{
_S85 = true;
}
else
{
if(j_0 == u32(0))
{
_S85 = (_S80.y) < u32(6);
}
else
{
_S85 = false;
}
}
if(_S85)
{
k_0 = k_0 + u32(1);
continue;
}
var _S86 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S83 + vec2<i32>(vec2<u32>(i_4, j_0))));
if((_S86.id_1) != u32(4294967295))
{
var _S87 : u32 = atomicMax(&(max_linked_list_length_0), entryPointParams_rigid_nodes_linked_lists_0[node_id_0(block_header_id_to_physical_id_0(_S86), _S80.xy).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( _S88 : vec3<u32>, _S89 : ptr<function, BlockVirtualId_std430_0>, _S90 : u32)
{
var _S91 : vec2<i32> = (*_S89).id_0 - vec2<i32>(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(0))
{
}
else
{
break;
}
var _S92 : bool;
if(i_5 == u32(0))
{
_S92 = (_S88.x) < u32(6);
}
else
{
_S92 = false;
}
var _S93 : bool;
if(_S92)
{
_S93 = true;
}
else
{
if(j_1 == u32(0))
{
_S93 = (_S88.y) < u32(6);
}
else
{
_S93 = false;
}
}
if(_S93)
{
k_1 = k_1 + u32(1);
continue;
}
var _S94 : vec2<u32> = vec2<u32>(i_5, j_1);
var _S95 : BlockHeaderId_0 = find_block_header_id_0(BlockVirtualId_x24init_0(_S91 + vec2<i32>(_S94)));
var _S96 : vec2<u32> = _S88.xy;
var _S97 : vec2<u32> = _S94 * vec2<u32>(u32(8)) + _S96;
var _S98 : u32 = flatten_shared_index_0(_S97.x, _S97.y);
if((_S95.id_1) != u32(4294967295))
{
var _S99 : NodePhysicalId_0 = node_id_0(block_header_id_to_physical_id_0(_S95), _S96);
shared_nodes_0[_S98].particle_id_0 = entryPointParams_rigid_nodes_linked_lists_0[_S99.id_3].head_0;
shared_nodes_0[_S98].global_id_0 = _S99.id_3;
}
else
{
shared_nodes_0[_S98].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( _S100 : 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(0))
{
}
else
{
break;
}
var _S101 : bool;
if(i_6 == u32(0))
{
_S101 = (_S100.x) < u32(6);
}
else
{
_S101 = false;
}
var _S102 : bool;
if(_S101)
{
_S102 = true;
}
else
{
if(j_2 == u32(0))
{
_S102 = (_S100.y) < u32(6);
}
else
{
_S102 = false;
}
}
if(_S102)
{
k_2 = k_2 + u32(1);
continue;
}
var _S103 : vec2<u32> = vec2<u32>(i_6, j_2) * vec2<u32>(u32(8)) + _S100.xy;
var _S104 : u32 = flatten_shared_index_0(_S103.x, _S103.y);
var _S105 : u32 = shared_nodes_0[_S104].particle_id_0;
if((shared_nodes_0[_S104].particle_id_0) != u32(4294967295))
{
var _S106 : RigidParticleIndices_std430_0 = entryPointParams_rigid_particle_indices_0[_S105];
shared_collider_ids_0[_S104] = _S106.collider_0;
shared_primitives_0[_S104] = Segment_x24init_0(entryPointParams_collider_vertices_0[_S106.segment_0.x], entryPointParams_collider_vertices_0[_S106.segment_0.y]);
shared_nodes_0[_S104].particle_id_0 = entryPointParams_particle_node_linked_lists_0[_S105];
}
else
{
shared_collider_ids_0[_S104] = u32(4294967295);
var _S107 : vec2<f32> = vec2<f32>(0.0f);
shared_primitives_0[_S104] = Segment_x24init_0(_S107, _S107);
}
k_2 = k_2 + u32(1);
}
j_2 = j_2 + u32(1);
}
i_6 = i_6 + u32(1);
}
return;
}
@compute
@workgroup_size(8, 8, 1)
fn p2g_cdf(@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 _S108 : u32 = block_id_0.x;
var _S109 : ActiveBlockHeaderGeneric_std430_0 = entryPointParams_active_blocks_0[_S108];
if(tid_flat_0 == u32(0))
{
atomicStore(&(max_linked_list_length_0), u32(0));
}
workgroupBarrier();
fetch_max_linked_lists_length_0(tid_0, &(_S109.virtual_id_0), _S108);
workgroupBarrier();
var _S110 : u32 = atomicLoad(&(max_linked_list_length_0));
max_linked_list_length_uniform_0 = _S110;
fetch_nodes_0(tid_0, &(_S109.virtual_id_0), _S108);
var _S111 : u32 = flatten_shared_index_0(tid_0.x + u32(8), tid_0.y + u32(8));
var _S112 : vec2<f32> = vec2<f32>(_S109.virtual_id_0.id_0 * vec2<i32>(i32(8)) + vec2<i32>(tid_0.xy)) * vec2<f32>(entryPointParams_grid_0[i32(0)].cell_width_0);
var _S113 : u32 = shared_nodes_0[_S111].global_id_0;
var node_cdf_0 : NodeCdf_0;
node_cdf_0.distance_0 = entryPointParams_nodes_0[shared_nodes_0[_S111].global_id_0].cdf_0.distance_0;
node_cdf_0.affinities_0 = entryPointParams_nodes_0[shared_nodes_0[_S111].global_id_0].cdf_0.affinities_0;
node_cdf_0.closest_id_0 = entryPointParams_nodes_0[shared_nodes_0[_S111].global_id_0].cdf_0.closest_id_0;
var _S114 : u32 = (workgroupUniformLoad(&((max_linked_list_length_uniform_0))));
var i_7 : u32 = u32(0);
for(;;)
{
if(i_7 < _S114)
{
}
else
{
break;
}
workgroupBarrier();
fetch_next_particle_0(tid_0);
workgroupBarrier();
var _S115 : NodeCdf_0 = p2g_step_0(_S111, entryPointParams_grid_0[i32(0)].cell_width_0, _S112);
if((_S115.closest_id_0) != u32(4294967295))
{
node_cdf_0.affinities_0 = ((node_cdf_0.affinities_0) | ((_S115.affinities_0)));
if((_S115.distance_0) < (node_cdf_0.distance_0))
{
node_cdf_0.distance_0 = _S115.distance_0;
node_cdf_0.closest_id_0 = _S115.closest_id_0;
}
}
i_7 = i_7 + u32(1);
}
entryPointParams_nodes_0[_S113].cdf_0.distance_0 = node_cdf_0.distance_0;
entryPointParams_nodes_0[_S113].cdf_0.affinities_0 = node_cdf_0.affinities_0;
entryPointParams_nodes_0[_S113].cdf_0.closest_id_0 = node_cdf_0.closest_id_0;
return;
}