archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

Fix #8151. (#8215)

Browse Source
This commit is contained in:
Cai Bear 2025-09-24 09:03:58 -07:00 committed by GitHub
parent ba0b6b9b0e
commit d2ee4b8be5
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 99 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
naga/src/front/spv/mod.rs
View File

@ -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)?;

63
naga/tests/in/spv/8151-barrier-reorder.spvasm Normal file
View File

@ -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

30
naga/tests/out/wgsl/spv-8151-barrier-reorder.wgsl Normal file
View File

@ -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();
}