Naga reorders workgroup barriers #8151
Description
Description
I tried translating the following SPIR-V to WGSL using naga and the workgroup barrier calls were reordered which breaks the shader.
Repro steps
let module = naga::front::spv::parse_u8_slice(&spirv, &Options::default()).unwrap();
let mut validator = Validator::new(ValidationFlags::all(), Capabilities::all());
let info = validator.validate(&module).unwrap();
let wgsl = naga::back::wgsl::write_string(&module, &info, WriterFlags::empty()).unwrap();; 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
Expected vs observed behavior
The resulting WGSL has the barriers in the wrong places such that the read from global_2 is not synchronized with the conditional write after it.
struct type_4 {
member: array<u32>,
}
var<private> global: vec3<u32>;
@group(0) @binding(0)
var<storage, read_write> global_1: type_4;
var<workgroup> global_2: u32;
fn function() {
let _e7 = global;
let _e9 = (_e7.x == 0u);
if _e9 {
global_2 = 1u;
}
workgroupBarrier();
workgroupBarrier();
let _e10 = global_2;
global_1.member[_e7.x] = _e10;
if _e9 {
global_2 = 2u;
}
return;
}
@compute @workgroup_size(2, 1, 1)
fn barrier_reorder_bug(@builtin(local_invocation_id) param: vec3<u32>) {
global = param;
function();
}Platform
naga = "26.0.0"