diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 13d0857f742..491416b664f 100644 --- a/naga/src/back/dot/mod.rs +++ b/naga/src/back/dot/mod.rs @@ -112,7 +112,8 @@ impl StatementGraph { } "Continue" } - S::Barrier(_flags) => "Barrier", + S::ControlBarrier(_flags) => "ControlBarrier", + S::MemoryBarrier(_flags) => "MemoryBarrier", S::Block(ref b) => { let (other, last) = self.add(b, targets); self.flow.push((id, other, "")); diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 81b740fc4f3..ef1621cfb90 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -1968,7 +1968,7 @@ impl<'a, W: Write> Writer<'a, W> { } writeln!(self.out, "{level}}}")?; - self.write_barrier(crate::Barrier::WORK_GROUP, level)?; + self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?; } Ok(()) @@ -2512,8 +2512,11 @@ impl<'a, W: Write> Writer<'a, W> { // keyword which ceases all further processing in a fragment shader, it's called OpKill // in spir-v that's why it's called `Statement::Kill` Statement::Kill => writeln!(self.out, "{level}discard;")?, - Statement::Barrier(flags) => { - self.write_barrier(flags, level)?; + Statement::ControlBarrier(flags) => { + self.write_control_barrier(flags, level)?; + } + Statement::MemoryBarrier(flags) => { + self.write_memory_barrier(flags, level)?; } // Stores in glsl are just variable assignments written as `pointer = value;` Statement::Store { pointer, value } => { @@ -2527,14 +2530,14 @@ impl<'a, W: Write> Writer<'a, W> { // GLSL doesn't have pointers, which means that this backend needs to ensure that // the actual "loading" is happening between the two barriers. // This is done in `Emit` by never emitting a variable name for pointer variables - self.write_barrier(crate::Barrier::WORK_GROUP, level)?; + self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?; let result_name = Baked(result).to_string(); write!(self.out, "{level}")?; // Expressions cannot have side effects, so just writing the expression here is fine. self.write_named_expr(pointer, result_name, result, ctx)?; - self.write_barrier(crate::Barrier::WORK_GROUP, level)?; + self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?; } // Stores a value into an image. Statement::ImageStore { @@ -4912,9 +4915,19 @@ impl<'a, W: Write> Writer<'a, W> { Ok(()) } - /// Issue a memory barrier. Please note that to ensure visibility, - /// OpenGL always requires a call to the `barrier()` function after a `memoryBarrier*()` - fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { + /// Issue a control barrier. + fn write_control_barrier( + &mut self, + flags: crate::Barrier, + level: back::Level, + ) -> BackendResult { + self.write_memory_barrier(flags, level)?; + writeln!(self.out, "{level}barrier();")?; + Ok(()) + } + + /// Issue a memory barrier. + fn write_memory_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { if flags.contains(crate::Barrier::STORAGE) { writeln!(self.out, "{level}memoryBarrierBuffer();")?; } @@ -4927,7 +4940,6 @@ impl<'a, W: Write> Writer<'a, W> { if flags.contains(crate::Barrier::TEXTURE) { writeln!(self.out, "{level}memoryBarrierImage();")?; } - writeln!(self.out, "{level}barrier();")?; Ok(()) } diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 192238d1f8d..8dc63a54469 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -1647,7 +1647,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } writeln!(self.out, "{level}}}")?; - self.write_barrier(crate::Barrier::WORK_GROUP, level) + self.write_control_barrier(crate::Barrier::WORK_GROUP, level) } /// Helper method used to write switches @@ -2291,8 +2291,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, "{level}continue;")? } } - Statement::Barrier(barrier) => { - self.write_barrier(barrier, level)?; + Statement::ControlBarrier(barrier) => { + self.write_control_barrier(barrier, level)?; + } + Statement::MemoryBarrier(barrier) => { + self.write_memory_barrier(barrier, level)?; } Statement::ImageStore { image, @@ -2464,12 +2467,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, ");")?; } Statement::WorkGroupUniformLoad { pointer, result } => { - self.write_barrier(crate::Barrier::WORK_GROUP, level)?; + self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?; write!(self.out, "{level}")?; let name = Baked(result).to_string(); self.write_named_expr(module, pointer, name, result, func_ctx)?; - self.write_barrier(crate::Barrier::WORK_GROUP, level)?; + self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?; } Statement::Switch { selector, @@ -4287,7 +4290,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Ok(()) } - fn write_barrier(&mut self, barrier: crate::Barrier, level: back::Level) -> BackendResult { + fn write_control_barrier( + &mut self, + barrier: crate::Barrier, + level: back::Level, + ) -> BackendResult { if barrier.contains(crate::Barrier::STORAGE) { writeln!(self.out, "{level}DeviceMemoryBarrierWithGroupSync();")?; } @@ -4303,6 +4310,26 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Ok(()) } + fn write_memory_barrier( + &mut self, + barrier: crate::Barrier, + level: back::Level, + ) -> BackendResult { + if barrier.contains(crate::Barrier::STORAGE) { + writeln!(self.out, "{level}DeviceMemoryBarrier();")?; + } + if barrier.contains(crate::Barrier::WORK_GROUP) { + writeln!(self.out, "{level}GroupMemoryBarrier();")?; + } + if barrier.contains(crate::Barrier::SUB_GROUP) { + // Does not exist in DirectX + } + if barrier.contains(crate::Barrier::TEXTURE) { + writeln!(self.out, "{level}DeviceMemoryBarrier();")?; + } + Ok(()) + } + /// Helper to emit the shared tail of an HLSL atomic call (arguments, value, result) fn emit_hlsl_atomic_tail( &mut self, diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index b9f0bd99479..899a0ffb415 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -3700,7 +3700,8 @@ impl Writer { crate::Statement::Kill => { writeln!(self.out, "{level}{NAMESPACE}::discard_fragment();")?; } - crate::Statement::Barrier(flags) => { + crate::Statement::ControlBarrier(flags) + | crate::Statement::MemoryBarrier(flags) => { self.write_barrier(flags, level)?; } crate::Statement::Store { pointer, value } => { diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 0088c6eac3f..62ce26de499 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -803,7 +803,11 @@ fn adjust_stmt(new_pos: &HandleVec>, stmt: &mut S crate::RayQueryFunction::Terminate => {} } } - Statement::Break | Statement::Continue | Statement::Kill | Statement::Barrier(_) => {} + Statement::Break + | Statement::Continue + | Statement::Kill + | Statement::ControlBarrier(_) + | Statement::MemoryBarrier(_) => {} } } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 8c1c8c4caa2..0dc3001bd0e 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -3240,8 +3240,11 @@ impl BlockContext<'_> { self.function.consume(block, Instruction::kill()); return Ok(BlockExitDisposition::Discarded); } - Statement::Barrier(flags) => { - self.writer.write_barrier(flags, &mut block); + Statement::ControlBarrier(flags) => { + self.writer.write_control_barrier(flags, &mut block); + } + Statement::MemoryBarrier(flags) => { + self.writer.write_memory_barrier(flags, &mut block); } Statement::Store { pointer, value } => { let value_id = self.cached[value]; @@ -3576,7 +3579,7 @@ impl BlockContext<'_> { } Statement::WorkGroupUniformLoad { pointer, result } => { self.writer - .write_barrier(crate::Barrier::WORK_GROUP, &mut block); + .write_control_barrier(crate::Barrier::WORK_GROUP, &mut block); let result_type_id = self.get_expression_type_id(&self.fun_info[result].ty); // Embed the body of match self.write_access_chain( @@ -3616,7 +3619,7 @@ impl BlockContext<'_> { } } self.writer - .write_barrier(crate::Barrier::WORK_GROUP, &mut block); + .write_control_barrier(crate::Barrier::WORK_GROUP, &mut block); } Statement::RayQuery { query, ref fun } => { self.write_ray_query_function(query, fun, &mut block); diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 97cf54587c6..788c3bc119a 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -1138,6 +1138,12 @@ impl super::Instruction { instruction.add_operand(semantics_id); instruction } + pub(super) fn memory_barrier(mem_scope_id: Word, semantics_id: Word) -> Self { + let mut instruction = Self::new(Op::MemoryBarrier); + instruction.add_operand(mem_scope_id); + instruction.add_operand(semantics_id); + instruction + } // Group Instructions diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 1fccec4a999..b0fd2068958 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -1711,9 +1711,11 @@ impl Writer { Ok(id) } - pub(super) fn write_barrier(&mut self, flags: crate::Barrier, block: &mut Block) { + pub(super) fn write_control_barrier(&mut self, flags: crate::Barrier, block: &mut Block) { let memory_scope = if flags.contains(crate::Barrier::STORAGE) { spirv::Scope::Device + } else if flags.contains(crate::Barrier::SUB_GROUP) { + spirv::Scope::Subgroup } else { spirv::Scope::Workgroup }; @@ -1726,6 +1728,10 @@ impl Writer { spirv::MemorySemantics::WORKGROUP_MEMORY, flags.contains(crate::Barrier::WORK_GROUP), ); + semantics.set( + spirv::MemorySemantics::SUBGROUP_MEMORY, + flags.contains(crate::Barrier::SUB_GROUP), + ); semantics.set( spirv::MemorySemantics::IMAGE_MEMORY, flags.contains(crate::Barrier::TEXTURE), @@ -1744,6 +1750,37 @@ impl Writer { )); } + pub(super) fn write_memory_barrier(&mut self, flags: crate::Barrier, block: &mut Block) { + let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE; + semantics.set( + spirv::MemorySemantics::UNIFORM_MEMORY, + flags.contains(crate::Barrier::STORAGE), + ); + semantics.set( + spirv::MemorySemantics::WORKGROUP_MEMORY, + flags.contains(crate::Barrier::WORK_GROUP), + ); + semantics.set( + spirv::MemorySemantics::SUBGROUP_MEMORY, + flags.contains(crate::Barrier::SUB_GROUP), + ); + semantics.set( + spirv::MemorySemantics::IMAGE_MEMORY, + flags.contains(crate::Barrier::TEXTURE), + ); + let mem_scope_id = if flags.contains(crate::Barrier::STORAGE) { + self.get_index_constant(spirv::Scope::Device as u32) + } else if flags.contains(crate::Barrier::SUB_GROUP) { + self.get_index_constant(spirv::Scope::Subgroup as u32) + } else { + self.get_index_constant(spirv::Scope::Workgroup as u32) + }; + let semantics_id = self.get_index_constant(semantics.bits()); + block + .body + .push(Instruction::memory_barrier(mem_scope_id, semantics_id)); + } + fn generate_workgroup_vars_init_block( &mut self, entry_id: Word, @@ -1844,7 +1881,7 @@ impl Writer { let mut post_if_block = Block::new(merge_id); - self.write_barrier(crate::Barrier::WORK_GROUP, &mut post_if_block); + self.write_control_barrier(crate::Barrier::WORK_GROUP, &mut post_if_block); let next_id = self.id_gen.next(); function.consume(post_if_block, Instruction::branch(next_id)); diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index 3ae12b3ecf1..e89fa3e425f 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -830,7 +830,7 @@ impl Writer { Statement::Continue => { writeln!(self.out, "{level}continue;")?; } - Statement::Barrier(barrier) => { + Statement::ControlBarrier(barrier) | Statement::MemoryBarrier(barrier) => { if barrier.contains(crate::Barrier::STORAGE) { writeln!(self.out, "{level}storageBarrier();")?; } diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index bf8d6ec7c8f..d7e7c4c6f22 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -155,7 +155,8 @@ impl FunctionTracer<'_> { St::Break | St::Continue | St::Kill - | St::Barrier(_) + | St::ControlBarrier(_) + | St::MemoryBarrier(_) | St::Return { value: None } => {} } } @@ -364,7 +365,8 @@ impl FunctionMap { St::Break | St::Continue | St::Kill - | St::Barrier(_) + | St::ControlBarrier(_) + | St::MemoryBarrier(_) | St::Return { value: None } => {} } } diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs index f73004aa9e8..37cf14fc4fe 100644 --- a/naga/src/front/glsl/builtins.rs +++ b/naga/src/front/glsl/builtins.rs @@ -2035,8 +2035,10 @@ impl MacroCall { )?, MacroCall::Barrier => { ctx.emit_restart(); - ctx.body - .push(crate::Statement::Barrier(crate::Barrier::all()), meta); + ctx.body.push( + crate::Statement::ControlBarrier(crate::Barrier::all()), + meta, + ); return Ok(None); } MacroCall::SmoothStep { splatted } => { diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index e59959d02ba..7e4f6c9be11 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -3850,7 +3850,9 @@ impl> Frontend { let semantics = resolve_constant(ctx.gctx(), &semantics_const.inner) .ok_or(Error::InvalidBarrierMemorySemantics(semantics_id))?; - if exec_scope == spirv::Scope::Workgroup as u32 { + if exec_scope == spirv::Scope::Workgroup as u32 + || exec_scope == spirv::Scope::Subgroup as u32 + { let mut flags = crate::Barrier::empty(); flags.set( crate::Barrier::STORAGE, @@ -3858,21 +3860,60 @@ impl> Frontend { ); flags.set( crate::Barrier::WORK_GROUP, - semantics - & (spirv::MemorySemantics::SUBGROUP_MEMORY - | spirv::MemorySemantics::WORKGROUP_MEMORY) - .bits() - != 0, + semantics & (spirv::MemorySemantics::WORKGROUP_MEMORY).bits() != 0, + ); + flags.set( + crate::Barrier::SUB_GROUP, + semantics & spirv::MemorySemantics::SUBGROUP_MEMORY.bits() != 0, ); flags.set( crate::Barrier::TEXTURE, semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0, ); - block.push(crate::Statement::Barrier(flags), span); + block.push(crate::Statement::ControlBarrier(flags), span); } else { log::warn!("Unsupported barrier execution scope: {}", exec_scope); } } + Op::MemoryBarrier => { + inst.expect(3)?; + let mem_scope_id = self.next()?; + let semantics_id = self.next()?; + let mem_scope_const = self.lookup_constant.lookup(mem_scope_id)?; + let semantics_const = self.lookup_constant.lookup(semantics_id)?; + + let mem_scope = resolve_constant(ctx.gctx(), &mem_scope_const.inner) + .ok_or(Error::InvalidBarrierScope(mem_scope_id))?; + let semantics = resolve_constant(ctx.gctx(), &semantics_const.inner) + .ok_or(Error::InvalidBarrierMemorySemantics(semantics_id))?; + + let mut flags = if mem_scope == spirv::Scope::Device as u32 { + crate::Barrier::STORAGE + } else if mem_scope == spirv::Scope::Workgroup as u32 { + crate::Barrier::WORK_GROUP + } else if mem_scope == spirv::Scope::Subgroup as u32 { + crate::Barrier::SUB_GROUP + } else { + crate::Barrier::empty() + }; + flags.set( + crate::Barrier::STORAGE, + semantics & spirv::MemorySemantics::UNIFORM_MEMORY.bits() != 0, + ); + flags.set( + crate::Barrier::WORK_GROUP, + semantics & (spirv::MemorySemantics::WORKGROUP_MEMORY).bits() != 0, + ); + flags.set( + crate::Barrier::SUB_GROUP, + semantics & spirv::MemorySemantics::SUBGROUP_MEMORY.bits() != 0, + ); + flags.set( + crate::Barrier::TEXTURE, + semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0, + ); + block.push(crate::Statement::MemoryBarrier(flags), span); + } Op::CopyObject => { inst.expect(4)?; let result_type_id = self.next()?; @@ -4566,7 +4607,8 @@ impl> Frontend { | S::Continue | S::Return { .. } | S::Kill - | S::Barrier(_) + | S::ControlBarrier(_) + | S::MemoryBarrier(_) | S::Store { .. } | S::ImageStore { .. } | S::Atomic { .. } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 4bf80018161..bd0f736974f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -2676,7 +2676,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let rctx = ctx.runtime_expression_ctx(span)?; rctx.block - .push(ir::Statement::Barrier(ir::Barrier::STORAGE), span); + .push(ir::Statement::ControlBarrier(ir::Barrier::STORAGE), span); return Ok(None); } "workgroupBarrier" => { @@ -2684,7 +2684,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let rctx = ctx.runtime_expression_ctx(span)?; rctx.block - .push(ir::Statement::Barrier(ir::Barrier::WORK_GROUP), span); + .push(ir::Statement::ControlBarrier(ir::Barrier::WORK_GROUP), span); return Ok(None); } "subgroupBarrier" => { @@ -2692,7 +2692,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let rctx = ctx.runtime_expression_ctx(span)?; rctx.block - .push(ir::Statement::Barrier(ir::Barrier::SUB_GROUP), span); + .push(ir::Statement::ControlBarrier(ir::Barrier::SUB_GROUP), span); return Ok(None); } "textureBarrier" => { @@ -2700,7 +2700,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let rctx = ctx.runtime_expression_ctx(span)?; rctx.block - .push(ir::Statement::Barrier(ir::Barrier::TEXTURE), span); + .push(ir::Statement::ControlBarrier(ir::Barrier::TEXTURE), span); return Ok(None); } "workgroupUniformLoad" => { diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs index 45478b8f8d1..14e88b761bd 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -1916,7 +1916,12 @@ pub enum Statement { /// Synchronize invocations within the work group. /// The `Barrier` flags control which memory accesses should be synchronized. /// If empty, this becomes purely an execution barrier. - Barrier(Barrier), + ControlBarrier(Barrier), + + /// Synchronize invocations within the work group. + /// The `Barrier` flags control which memory accesses should be synchronized. + MemoryBarrier(Barrier), + /// Stores a value at an address. /// /// For [`TypeInner::Atomic`] type behind the pointer, the value diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs index f22e61e6a6d..b29ccb054a3 100644 --- a/naga/src/proc/terminator.rs +++ b/naga/src/proc/terminator.rs @@ -42,7 +42,8 @@ pub fn ensure_block_returns(block: &mut crate::Block) { | S::SubgroupBallot { .. } | S::SubgroupCollectiveOperation { .. } | S::SubgroupGather { .. } - | S::Barrier(_)), + | S::ControlBarrier(_) + | S::MemoryBarrier(_)), ) | None => block.push(S::Return { value: None }, Default::default()), } diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index 435e6b9fd57..857dfdc14c1 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -902,7 +902,7 @@ impl FunctionInfo { ExitFlags::empty() }, }, - S::Barrier(_) => FunctionUniformity { + S::ControlBarrier(_) | S::MemoryBarrier(_) => FunctionUniformity { result: Uniformity { non_uniform_result: None, requirements: UniformityRequirements::WORK_GROUP_BARRIER, diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 513a92aa164..fb88f3934f5 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1006,7 +1006,7 @@ impl super::Validator { S::Kill => { stages &= super::ShaderStages::FRAGMENT; } - S::Barrier(barrier) => { + S::ControlBarrier(barrier) | S::MemoryBarrier(barrier) => { stages &= super::ShaderStages::COMPUTE; if barrier.contains(crate::Barrier::SUB_GROUP) { if !self.capabilities.contains( diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs index 9f59edc0ec5..15cbfe008e7 100644 --- a/naga/src/valid/handles.rs +++ b/naga/src/valid/handles.rs @@ -837,7 +837,8 @@ impl super::Validator { crate::Statement::Break | crate::Statement::Continue | crate::Statement::Kill - | crate::Statement::Barrier(_) => Ok(()), + | crate::Statement::ControlBarrier(_) + | crate::Statement::MemoryBarrier(_) => Ok(()), }) } } diff --git a/naga/tests/in/spv/barrier.spvasm b/naga/tests/in/spv/barrier.spvasm new file mode 100644 index 00000000000..bc1341c8262 --- /dev/null +++ b/naga/tests/in/spv/barrier.spvasm @@ -0,0 +1,27 @@ +; SPIR-V +; Version: 1.5 +; Generator: Google rspirv; 0 +; Bound: 14 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical Simple + OpEntryPoint GLCompute %1 "main" + OpExecutionMode %1 LocalSize 64 1 1 + %void = OpTypeVoid + %6 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 + %uint_1 = OpConstant %uint 1 + %uint_2120 = OpConstant %uint 2120 + %uint_2376 = OpConstant %uint 2376 + %1 = OpFunction %void None %6 + %13 = OpLabel + OpMemoryBarrier %uint_2 %uint_264 + OpControlBarrier %uint_2 %uint_2 %uint_264 + OpMemoryBarrier %uint_1 %uint_2120 + OpControlBarrier %uint_2 %uint_1 %uint_2120 + OpMemoryBarrier %uint_1 %uint_2376 + OpControlBarrier %uint_2 %uint_1 %uint_2376 + OpReturn + OpFunctionEnd diff --git a/naga/tests/in/spv/barrier.toml b/naga/tests/in/spv/barrier.toml new file mode 100644 index 00000000000..6a727192212 --- /dev/null +++ b/naga/tests/in/spv/barrier.toml @@ -0,0 +1,4 @@ +targets = "WGSL | SPIRV | GLSL | HLSL | METAL" + +[msl] +lang_version = [2, 0] diff --git a/naga/tests/in/spv/subgroup-barrier.spvasm b/naga/tests/in/spv/subgroup-barrier.spvasm new file mode 100644 index 00000000000..c3f97cff2e2 --- /dev/null +++ b/naga/tests/in/spv/subgroup-barrier.spvasm @@ -0,0 +1,20 @@ +; SPIR-V +; Version: 1.5 +; Generator: Google rspirv; 0 +; Bound: 14 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical Simple + OpEntryPoint GLCompute %1 "main" + OpExecutionMode %1 LocalSize 64 1 1 + %void = OpTypeVoid + %6 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_3 = OpConstant %uint 3 + %uint_136 = OpConstant %uint 136 + %1 = OpFunction %void None %6 + %13 = OpLabel + OpMemoryBarrier %uint_3 %uint_136 + OpControlBarrier %uint_3 %uint_3 %uint_136 + OpReturn + OpFunctionEnd diff --git a/naga/tests/in/spv/subgroup-barrier.toml b/naga/tests/in/spv/subgroup-barrier.toml new file mode 100644 index 00000000000..48ca9914b31 --- /dev/null +++ b/naga/tests/in/spv/subgroup-barrier.toml @@ -0,0 +1,5 @@ +god_mode = true +targets = "WGSL | SPIRV | GLSL | METAL" + +[msl] +lang_version = [2, 0] diff --git a/naga/tests/out/glsl/spv-barrier.main.Compute.glsl b/naga/tests/out/glsl/spv-barrier.main.Compute.glsl new file mode 100644 index 00000000000..a5be41cf0b1 --- /dev/null +++ b/naga/tests/out/glsl/spv-barrier.main.Compute.glsl @@ -0,0 +1,31 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in; + + +void function() { + memoryBarrierShared(); + memoryBarrierShared(); + barrier(); + memoryBarrierBuffer(); + memoryBarrierImage(); + memoryBarrierBuffer(); + memoryBarrierImage(); + barrier(); + memoryBarrierBuffer(); + memoryBarrierShared(); + memoryBarrierImage(); + memoryBarrierBuffer(); + memoryBarrierShared(); + memoryBarrierImage(); + barrier(); + return; +} + +void main() { + function(); +} + diff --git a/naga/tests/out/glsl/spv-subgroup-barrier.main.Compute.glsl b/naga/tests/out/glsl/spv-subgroup-barrier.main.Compute.glsl new file mode 100644 index 00000000000..ae961986737 --- /dev/null +++ b/naga/tests/out/glsl/spv-subgroup-barrier.main.Compute.glsl @@ -0,0 +1,19 @@ +#version 310 es + +precision highp float; +precision highp int; + +layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in; + + +void function() { + subgroupMemoryBarrier(); + subgroupMemoryBarrier(); + barrier(); + return; +} + +void main() { + function(); +} + diff --git a/naga/tests/out/glsl/spv-subgroup-operations-s.main.Compute.glsl b/naga/tests/out/glsl/spv-subgroup-operations-s.main.Compute.glsl index 67389282e65..1e591d8114d 100644 --- a/naga/tests/out/glsl/spv-subgroup-operations-s.main.Compute.glsl +++ b/naga/tests/out/glsl/spv-subgroup-operations-s.main.Compute.glsl @@ -21,6 +21,7 @@ uint global_3 = 0u; void function() { uint _e5 = global_2; uint _e6 = global_3; + barrier(); uvec4 _e9 = subgroupBallot(((_e6 & 1u) == 1u)); uvec4 _e10 = subgroupBallot(true); bool _e12 = subgroupAll((_e6 != 0u)); diff --git a/naga/tests/out/hlsl/spv-barrier.hlsl b/naga/tests/out/hlsl/spv-barrier.hlsl new file mode 100644 index 00000000000..e808dc9f990 --- /dev/null +++ b/naga/tests/out/hlsl/spv-barrier.hlsl @@ -0,0 +1,22 @@ +void function() +{ + GroupMemoryBarrier(); + GroupMemoryBarrierWithGroupSync(); + DeviceMemoryBarrier(); + DeviceMemoryBarrier(); + DeviceMemoryBarrierWithGroupSync(); + DeviceMemoryBarrierWithGroupSync(); + DeviceMemoryBarrier(); + GroupMemoryBarrier(); + DeviceMemoryBarrier(); + DeviceMemoryBarrierWithGroupSync(); + GroupMemoryBarrierWithGroupSync(); + DeviceMemoryBarrierWithGroupSync(); + return; +} + +[numthreads(64, 1, 1)] +void main() +{ + function(); +} diff --git a/naga/tests/out/hlsl/spv-barrier.ron b/naga/tests/out/hlsl/spv-barrier.ron new file mode 100644 index 00000000000..a07b03300b1 --- /dev/null +++ b/naga/tests/out/hlsl/spv-barrier.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/msl/spv-barrier.msl b/naga/tests/out/msl/spv-barrier.msl new file mode 100644 index 00000000000..faf14bfb7d4 --- /dev/null +++ b/naga/tests/out/msl/spv-barrier.msl @@ -0,0 +1,28 @@ +// language: metal2.0 +#include +#include + +using metal::uint; + + +void function( +) { + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::threadgroup_barrier(metal::mem_flags::mem_device); + metal::threadgroup_barrier(metal::mem_flags::mem_texture); + metal::threadgroup_barrier(metal::mem_flags::mem_device); + metal::threadgroup_barrier(metal::mem_flags::mem_texture); + metal::threadgroup_barrier(metal::mem_flags::mem_device); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::threadgroup_barrier(metal::mem_flags::mem_texture); + metal::threadgroup_barrier(metal::mem_flags::mem_device); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::threadgroup_barrier(metal::mem_flags::mem_texture); + return; +} + +kernel void main_( +) { + function(); +} diff --git a/naga/tests/out/msl/spv-subgroup-barrier.msl b/naga/tests/out/msl/spv-subgroup-barrier.msl new file mode 100644 index 00000000000..ce56e7a7ea1 --- /dev/null +++ b/naga/tests/out/msl/spv-subgroup-barrier.msl @@ -0,0 +1,18 @@ +// language: metal2.0 +#include +#include + +using metal::uint; + + +void function( +) { + metal::simdgroup_barrier(metal::mem_flags::mem_threadgroup); + metal::simdgroup_barrier(metal::mem_flags::mem_threadgroup); + return; +} + +kernel void main_( +) { + function(); +} diff --git a/naga/tests/out/msl/spv-subgroup-operations-s.msl b/naga/tests/out/msl/spv-subgroup-operations-s.msl index 163dd46649e..fcd42730376 100644 --- a/naga/tests/out/msl/spv-subgroup-operations-s.msl +++ b/naga/tests/out/msl/spv-subgroup-operations-s.msl @@ -11,6 +11,7 @@ void function( ) { uint _e5 = global_2; uint _e6 = global_3; + metal::threadgroup_barrier(metal::mem_flags::mem_none); metal::uint4 unnamed = metal::uint4((uint64_t)metal::simd_ballot((_e6 & 1u) == 1u), 0, 0, 0); metal::uint4 unnamed_1 = metal::uint4((uint64_t)metal::simd_ballot(true), 0, 0, 0); bool unnamed_2 = metal::simd_all(_e6 != 0u); diff --git a/naga/tests/out/spv/spv-barrier.spvasm b/naga/tests/out/spv/spv-barrier.spvasm new file mode 100644 index 00000000000..9a5d61cf93d --- /dev/null +++ b/naga/tests/out/spv/spv-barrier.spvasm @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 17 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %14 "main" +OpExecutionMode %14 LocalSize 64 1 1 +%2 = OpTypeVoid +%5 = OpTypeFunction %2 +%8 = OpTypeInt 32 0 +%7 = OpConstant %8 2 +%9 = OpConstant %8 264 +%10 = OpConstant %8 1 +%11 = OpConstant %8 2120 +%12 = OpConstant %8 2376 +%4 = OpFunction %2 None %5 +%3 = OpLabel +OpBranch %6 +%6 = OpLabel +OpMemoryBarrier %7 %9 +OpControlBarrier %7 %7 %9 +OpMemoryBarrier %10 %11 +OpControlBarrier %7 %10 %11 +OpMemoryBarrier %10 %12 +OpControlBarrier %7 %10 %12 +OpReturn +OpFunctionEnd +%14 = OpFunction %2 None %5 +%13 = OpLabel +OpBranch %15 +%15 = OpLabel +%16 = OpFunctionCall %2 %4 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/spv-subgroup-barrier.spvasm b/naga/tests/out/spv/spv-subgroup-barrier.spvasm new file mode 100644 index 00000000000..8aac1b39f31 --- /dev/null +++ b/naga/tests/out/spv/spv-subgroup-barrier.spvasm @@ -0,0 +1,29 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 14 +OpCapability Shader +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %11 "main" +OpExecutionMode %11 LocalSize 64 1 1 +%2 = OpTypeVoid +%5 = OpTypeFunction %2 +%8 = OpTypeInt 32 0 +%7 = OpConstant %8 3 +%9 = OpConstant %8 136 +%4 = OpFunction %2 None %5 +%3 = OpLabel +OpBranch %6 +%6 = OpLabel +OpMemoryBarrier %7 %9 +OpControlBarrier %7 %7 %9 +OpReturn +OpFunctionEnd +%11 = OpFunction %2 None %5 +%10 = OpLabel +OpBranch %12 +%12 = OpLabel +%13 = OpFunctionCall %2 %4 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-subgroup-operations.spvasm b/naga/tests/out/spv/wgsl-subgroup-operations.spvasm index f4f257ad0ec..b36eb95a852 100644 --- a/naga/tests/out/spv/wgsl-subgroup-operations.spvasm +++ b/naga/tests/out/spv/wgsl-subgroup-operations.spvasm @@ -34,10 +34,10 @@ OpDecorate %15 BuiltIn SubgroupLocalInvocationId %20 = OpConstant %3 0 %21 = OpConstant %3 4 %23 = OpConstant %3 3 -%24 = OpConstant %3 2 -%25 = OpConstant %3 8 -%28 = OpTypeVector %3 4 -%30 = OpConstantTrue %5 +%24 = OpConstant %3 136 +%27 = OpTypeVector %3 4 +%29 = OpConstantTrue %5 +%61 = OpConstant %3 2 %17 = OpFunction %2 None %18 %6 = OpLabel %10 = OpLoad %3 %8 @@ -47,40 +47,40 @@ OpDecorate %15 BuiltIn SubgroupLocalInvocationId %16 = OpLoad %3 %15 OpBranch %22 %22 = OpLabel -OpControlBarrier %23 %24 %25 -%26 = OpBitwiseAnd %3 %16 %19 -%27 = OpIEqual %5 %26 %19 -%29 = OpGroupNonUniformBallot %28 %23 %27 -%31 = OpGroupNonUniformBallot %28 %23 %30 -%32 = OpINotEqual %5 %16 %20 -%33 = OpGroupNonUniformAll %5 %23 %32 -%34 = OpIEqual %5 %16 %20 -%35 = OpGroupNonUniformAny %5 %23 %34 -%36 = OpGroupNonUniformIAdd %3 %23 Reduce %16 -%37 = OpGroupNonUniformIMul %3 %23 Reduce %16 -%38 = OpGroupNonUniformUMin %3 %23 Reduce %16 -%39 = OpGroupNonUniformUMax %3 %23 Reduce %16 -%40 = OpGroupNonUniformBitwiseAnd %3 %23 Reduce %16 -%41 = OpGroupNonUniformBitwiseOr %3 %23 Reduce %16 -%42 = OpGroupNonUniformBitwiseXor %3 %23 Reduce %16 -%43 = OpGroupNonUniformIAdd %3 %23 ExclusiveScan %16 -%44 = OpGroupNonUniformIMul %3 %23 ExclusiveScan %16 -%45 = OpGroupNonUniformIAdd %3 %23 InclusiveScan %16 -%46 = OpGroupNonUniformIMul %3 %23 InclusiveScan %16 -%47 = OpGroupNonUniformBroadcastFirst %3 %23 %16 -%48 = OpGroupNonUniformShuffle %3 %23 %16 %21 -%49 = OpCompositeExtract %3 %7 1 -%50 = OpISub %3 %49 %19 -%51 = OpISub %3 %50 %16 -%52 = OpGroupNonUniformShuffle %3 %23 %16 %51 -%53 = OpGroupNonUniformShuffleDown %3 %23 %16 %19 -%54 = OpGroupNonUniformShuffleUp %3 %23 %16 %19 -%55 = OpCompositeExtract %3 %7 1 -%56 = OpISub %3 %55 %19 -%57 = OpGroupNonUniformShuffleXor %3 %23 %16 %56 -%58 = OpGroupNonUniformQuadBroadcast %3 %23 %16 %21 -%59 = OpGroupNonUniformQuadSwap %3 %23 %16 %20 -%60 = OpGroupNonUniformQuadSwap %3 %23 %16 %19 -%61 = OpGroupNonUniformQuadSwap %3 %23 %16 %24 +OpControlBarrier %23 %23 %24 +%25 = OpBitwiseAnd %3 %16 %19 +%26 = OpIEqual %5 %25 %19 +%28 = OpGroupNonUniformBallot %27 %23 %26 +%30 = OpGroupNonUniformBallot %27 %23 %29 +%31 = OpINotEqual %5 %16 %20 +%32 = OpGroupNonUniformAll %5 %23 %31 +%33 = OpIEqual %5 %16 %20 +%34 = OpGroupNonUniformAny %5 %23 %33 +%35 = OpGroupNonUniformIAdd %3 %23 Reduce %16 +%36 = OpGroupNonUniformIMul %3 %23 Reduce %16 +%37 = OpGroupNonUniformUMin %3 %23 Reduce %16 +%38 = OpGroupNonUniformUMax %3 %23 Reduce %16 +%39 = OpGroupNonUniformBitwiseAnd %3 %23 Reduce %16 +%40 = OpGroupNonUniformBitwiseOr %3 %23 Reduce %16 +%41 = OpGroupNonUniformBitwiseXor %3 %23 Reduce %16 +%42 = OpGroupNonUniformIAdd %3 %23 ExclusiveScan %16 +%43 = OpGroupNonUniformIMul %3 %23 ExclusiveScan %16 +%44 = OpGroupNonUniformIAdd %3 %23 InclusiveScan %16 +%45 = OpGroupNonUniformIMul %3 %23 InclusiveScan %16 +%46 = OpGroupNonUniformBroadcastFirst %3 %23 %16 +%47 = OpGroupNonUniformShuffle %3 %23 %16 %21 +%48 = OpCompositeExtract %3 %7 1 +%49 = OpISub %3 %48 %19 +%50 = OpISub %3 %49 %16 +%51 = OpGroupNonUniformShuffle %3 %23 %16 %50 +%52 = OpGroupNonUniformShuffleDown %3 %23 %16 %19 +%53 = OpGroupNonUniformShuffleUp %3 %23 %16 %19 +%54 = OpCompositeExtract %3 %7 1 +%55 = OpISub %3 %54 %19 +%56 = OpGroupNonUniformShuffleXor %3 %23 %16 %55 +%57 = OpGroupNonUniformQuadBroadcast %3 %23 %16 %21 +%58 = OpGroupNonUniformQuadSwap %3 %23 %16 %20 +%59 = OpGroupNonUniformQuadSwap %3 %23 %16 %19 +%60 = OpGroupNonUniformQuadSwap %3 %23 %16 %61 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/spv-barrier.wgsl b/naga/tests/out/wgsl/spv-barrier.wgsl new file mode 100644 index 00000000000..dc8ad2cbff3 --- /dev/null +++ b/naga/tests/out/wgsl/spv-barrier.wgsl @@ -0,0 +1,20 @@ +fn function() { + workgroupBarrier(); + workgroupBarrier(); + storageBarrier(); + textureBarrier(); + storageBarrier(); + textureBarrier(); + storageBarrier(); + workgroupBarrier(); + textureBarrier(); + storageBarrier(); + workgroupBarrier(); + textureBarrier(); + return; +} + +@compute @workgroup_size(64, 1, 1) +fn main() { + function(); +} diff --git a/naga/tests/out/wgsl/spv-subgroup-barrier.wgsl b/naga/tests/out/wgsl/spv-subgroup-barrier.wgsl new file mode 100644 index 00000000000..a26a73e8c3f --- /dev/null +++ b/naga/tests/out/wgsl/spv-subgroup-barrier.wgsl @@ -0,0 +1,10 @@ +fn function() { + subgroupBarrier(); + subgroupBarrier(); + return; +} + +@compute @workgroup_size(64, 1, 1) +fn main() { + function(); +}