#version 460
#extension GL_KHR_shader_subgroup_ballot : enable
#extension GL_KHR_shader_subgroup_vote : enable
#extension GL_KHR_shader_subgroup_basic : enable
#extension GL_KHR_memory_scope_semantics : enable
layout(set = 0, binding = 0) buffer Buf {
 uint buf[];
};
layout(local_size_x = 128, local_size_y = 1, local_size_z = 1) in;
void main()
{
 uint subgroup_id = gl_SubgroupID;
uint subgroup_size = 16;
 uint subgroup_local_id = gl_SubgroupInvocationID;
 uint num_workgroup = gl_NumWorkGroups.x;
 uint workgroup_size = gl_WorkGroupSize.x;
 uint workgroup_id = gl_WorkGroupID.x;
 uint workgroup_base = workgroup_size * workgroup_id;
 uint virtual_gid = workgroup_base + subgroup_id * subgroup_size + subgroup_local_id;
 uint next_virtual_gid = workgroup_base + subgroup_id * subgroup_size + ((subgroup_local_id + 1) % subgroup_size);
 buf[virtual_gid] = 1;
 atomicStore(buf[next_virtual_gid], uint(2), 4, 64, 4);
}