Fix #8151, SPIR-V frontend reordering barriers

naga/src/front/spv/mod.rs

@@ -3869,7 +3869,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                             crate::Barrier::TEXTURE,
                             semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0,
                         );
+
+                        block.extend(emitter.finish(ctx.expressions));
                         block.push(crate::Statement::ControlBarrier(flags), span);
+                        emitter.start(ctx.expressions);
                     } else {
                         log::warn!("Unsupported barrier execution scope: {exec_scope}");
                     }
@@ -3911,7 +3914,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                         crate::Barrier::TEXTURE,
                         semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0,
                     );
+
+                    block.extend(emitter.finish(ctx.expressions));
                     block.push(crate::Statement::MemoryBarrier(flags), span);
+                    emitter.start(ctx.expressions);
                 }
                 Op::CopyObject => {
                     inst.expect(4)?;

naga/tests/in/spv/8151-barrier-reorder.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google rspirv; 0
+; Bound: 45
+; Schema: 0
+               OpCapability Shader
+               OpCapability VulkanMemoryModel
+               OpExtension "SPV_KHR_vulkan_memory_model"
+               OpMemoryModel Logical Vulkan
+               OpEntryPoint GLCompute %1 "barrier_reorder_bug" %gl_LocalInvocationID
+               OpExecutionMode %1 LocalSize 2 1 1
+               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpDecorate %_struct_8 Block
+               OpMemberDecorate %_struct_8 0 Offset 0
+               OpDecorate %5 Binding 0
+               OpDecorate %5 DescriptorSet 0
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+  %_struct_8 = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer__struct_8 = OpTypePointer StorageBuffer %_struct_8
+          %5 = OpVariable %_ptr_StorageBuffer__struct_8 StorageBuffer
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+          %6 = OpVariable %_ptr_Workgroup_uint Workgroup
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+          %1 = OpFunction %void None %13
+         %23 = OpLabel
+         %24 = OpLoad %v3uint %gl_LocalInvocationID
+         %27 = OpCompositeExtract %uint %24 0
+         %28 = OpIEqual %bool %27 %uint_0
+               OpSelectionMerge %29 None
+               OpBranchConditional %28 %30 %31
+         %30 = OpLabel
+               OpStore %6 %uint_1
+               OpBranch %29
+         %31 = OpLabel
+               OpBranch %29
+         %29 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %32 = OpLoad %uint %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %39 = OpInBoundsAccessChain %_ptr_StorageBuffer_uint %5 %uint_0 %27
+               OpStore %39 %32
+               OpSelectionMerge %42 None
+               OpBranchConditional %28 %43 %44
+         %43 = OpLabel
+               OpStore %6 %uint_2
+               OpBranch %42
+         %44 = OpLabel
+               OpBranch %42
+         %42 = OpLabel
+               OpReturn
+               OpFunctionEnd

naga/tests/out/wgsl/spv-8151-barrier-reorder.wgsl

@@ -0,0 +1,30 @@
+struct type_3 {
+    member: array<u32>,
+}
+
+var<private> global: vec3<u32>;
+@group(0) @binding(0) 
+var<storage, read_write> global_1: type_3;
+var<workgroup> global_2: u32;
+
+fn function() {
+    let _e6 = global;
+    let _e8 = (_e6.x == 0u);
+    if _e8 {
+        global_2 = 1u;
+    }
+    workgroupBarrier();
+    let _e9 = global_2;
+    workgroupBarrier();
+    global_1.member[_e6.x] = _e9;
+    if _e8 {
+        global_2 = 2u;
+    }
+    return;
+}
+
+@compute @workgroup_size(2, 1, 1) 
+fn barrier_reorder_bug(@builtin(local_invocation_id) param: vec3<u32>) {
+    global = param;
+    function();
+}