diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index 960437ece58..2e54c74a213 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -3869,7 +3869,10 @@ impl> Frontend { 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> Frontend { 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)?; diff --git a/naga/tests/in/spv/8151-barrier-reorder.spvasm b/naga/tests/in/spv/8151-barrier-reorder.spvasm new file mode 100644 index 00000000000..733e4d74414 --- /dev/null +++ b/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 diff --git a/naga/tests/out/wgsl/spv-8151-barrier-reorder.wgsl b/naga/tests/out/wgsl/spv-8151-barrier-reorder.wgsl new file mode 100644 index 00000000000..4ba429bb811 --- /dev/null +++ b/naga/tests/out/wgsl/spv-8151-barrier-reorder.wgsl @@ -0,0 +1,30 @@ +struct type_3 { + member: array, +} + +var global: vec3; +@group(0) @binding(0) +var global_1: type_3; +var 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) { + global = param; + function(); +}