diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index b3075c9e7c..b3d2682033 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -71,7 +71,7 @@ pub type BindingMap = std::collections::BTreeMap; impl crate::AtomicFunction { fn to_glsl(self) -> &'static str { match self { - Self::Add => "Add", + Self::Add | Self::Subtract => "Add", Self::And => "And", Self::InclusiveOr => "Or", Self::ExclusiveOr => "Xor", @@ -1706,12 +1706,20 @@ impl<'a, W: Write> Writer<'a, W> { let fun_str = fun.to_glsl(); write!(self.out, "atomic{}(", fun_str)?; self.write_expr(pointer, ctx)?; - if let crate::AtomicFunction::Exchange { compare: Some(_) } = *fun { - return Err(Error::Custom( - "atomic CompareExchange is not implemented".to_string(), - )); - } write!(self.out, ", ")?; + // handle the special cases + match *fun { + crate::AtomicFunction::Subtract => { + // we just wrote `InterlockedAdd`, so negate the argument + write!(self.out, "-")?; + } + crate::AtomicFunction::Exchange { compare: Some(_) } => { + return Err(Error::Custom( + "atomic CompareExchange is not implemented".to_string(), + )); + } + _ => {} + } self.write_expr(value, ctx)?; writeln!(self.out, ");")?; } diff --git a/src/back/hlsl/conv.rs b/src/back/hlsl/conv.rs index e535d45343..b8cb10daa2 100644 --- a/src/back/hlsl/conv.rs +++ b/src/back/hlsl/conv.rs @@ -133,7 +133,7 @@ impl crate::AtomicFunction { /// Return the HLSL suffix for the `InterlockedXxx` method. pub(super) fn to_hlsl_suffix(self) -> &'static str { match self { - Self::Add => "Add", + Self::Add | Self::Subtract => "Add", Self::And => "And", Self::InclusiveOr => "Or", Self::ExclusiveOr => "Xor", diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 96740fbec4..136aee3ca4 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1374,10 +1374,18 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { res_name, var_name, fun_str )?; self.write_storage_address(module, &chain, func_ctx)?; - if let crate::AtomicFunction::Exchange { compare: Some(_) } = *fun { - return Err(Error::Unimplemented("atomic CompareExchange".to_string())); - } write!(self.out, ", ")?; + // handle the special cases + match *fun { + crate::AtomicFunction::Subtract => { + // we just wrote `InterlockedAdd`, so negate the argument + write!(self.out, "-")?; + } + crate::AtomicFunction::Exchange { compare: Some(_) } => { + return Err(Error::Unimplemented("atomic CompareExchange".to_string())); + } + _ => {} + } self.write_expr(module, value, func_ctx)?; writeln!(self.out, ", {});", res_name)?; self.temp_access_chain = chain; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index a1be83b742..aa65970557 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1647,6 +1647,9 @@ impl Writer { crate::AtomicFunction::Add => { self.put_atomic_fetch(pointer, "add", value, &context.expression)?; } + crate::AtomicFunction::Subtract => { + self.put_atomic_fetch(pointer, "sub", value, &context.expression)?; + } crate::AtomicFunction::And => { self.put_atomic_fetch(pointer, "and", value, &context.expression)?; } @@ -2643,8 +2646,8 @@ fn test_stack_size() { } let stack_size = addresses.end - addresses.start; // check the size (in debug only) - // last observed macOS value: 17504 - if !(13000..=19000).contains(&stack_size) { + // last observed macOS value: 19152 (CI) + if !(13000..=20000).contains(&stack_size) { panic!("`put_block` stack size {} has changed!", stack_size); } } diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index f1e95f7719..3009085cee 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -1262,6 +1262,15 @@ impl<'w> BlockContext<'w> { semantics_id, value_id, ), + crate::AtomicFunction::Subtract => Instruction::atomic_binary( + spirv::Op::AtomicISub, + result_type_id, + id, + pointer_id, + scope_constant_id, + semantics_id, + value_id, + ), crate::AtomicFunction::And => Instruction::atomic_binary( spirv::Op::AtomicAnd, result_type_id, diff --git a/src/back/wgsl/mod.rs b/src/back/wgsl/mod.rs index 33d01e2458..c99df6d811 100644 --- a/src/back/wgsl/mod.rs +++ b/src/back/wgsl/mod.rs @@ -31,6 +31,7 @@ impl crate::AtomicFunction { fn to_wgsl(self) -> &'static str { match self { Self::Add => "Add", + Self::Subtract => "Sub", Self::And => "And", Self::InclusiveOr => "Or", Self::ExclusiveOr => "Xor", diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index bf97c8e0c7..ba3f54a98f 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -1403,6 +1403,15 @@ impl Parser { )?; return Ok(Some(handle)); } + "atomicSub" => { + let _ = lexer.next(); + let handle = self.parse_atomic_helper( + lexer, + crate::AtomicFunction::Subtract, + ctx.reborrow(), + )?; + return Ok(Some(handle)); + } "atomicAnd" => { let _ = lexer.next(); let handle = self.parse_atomic_helper( diff --git a/src/lib.rs b/src/lib.rs index 0b7f02ddd2..f672604561 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -750,6 +750,7 @@ pub enum BinaryOperator { #[cfg_attr(feature = "deserialize", derive(Deserialize))] pub enum AtomicFunction { Add, + Subtract, And, ExclusiveOr, InclusiveOr, diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index ba54363ed1..8ab506c276 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -43,6 +43,7 @@ fn atomics() { var tmp: i32; let value = atomicLoad(&bar.atom); tmp = atomicAdd(&bar.atom, 5); + tmp = atomicSub(&bar.atom, 5); tmp = atomicAnd(&bar.atom, 5); tmp = atomicOr(&bar.atom, 5); tmp = atomicXor(&bar.atom, 5); diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index f0a982862d..2af1e78275 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -18,18 +18,20 @@ void main() { int value = _group_0_binding_0.atom; int _e6 = atomicAdd(_group_0_binding_0.atom, 5); tmp = _e6; - int _e9 = atomicAnd(_group_0_binding_0.atom, 5); + int _e9 = atomicAdd(_group_0_binding_0.atom, -5); tmp = _e9; - int _e12 = atomicOr(_group_0_binding_0.atom, 5); + int _e12 = atomicAnd(_group_0_binding_0.atom, 5); tmp = _e12; - int _e15 = atomicXor(_group_0_binding_0.atom, 5); + int _e15 = atomicOr(_group_0_binding_0.atom, 5); tmp = _e15; - int _e18 = atomicMin(_group_0_binding_0.atom, 5); + int _e18 = atomicXor(_group_0_binding_0.atom, 5); tmp = _e18; - int _e21 = atomicMax(_group_0_binding_0.atom, 5); + int _e21 = atomicMin(_group_0_binding_0.atom, 5); tmp = _e21; - int _e24 = atomicExchange(_group_0_binding_0.atom, 5); + int _e24 = atomicMax(_group_0_binding_0.atom, 5); tmp = _e24; + int _e27 = atomicExchange(_group_0_binding_0.atom, 5); + tmp = _e27; _group_0_binding_0.atom = value; return; } diff --git a/tests/out/hlsl/access.hlsl b/tests/out/hlsl/access.hlsl index 9e306fcd8f..d9c765e03b 100644 --- a/tests/out/hlsl/access.hlsl +++ b/tests/out/hlsl/access.hlsl @@ -50,18 +50,20 @@ void atomics() int value = asint(bar.Load(64)); int _e6; bar.InterlockedAdd(64, 5, _e6); tmp = _e6; - int _e9; bar.InterlockedAnd(64, 5, _e9); + int _e9; bar.InterlockedAdd(64, -5, _e9); tmp = _e9; - int _e12; bar.InterlockedOr(64, 5, _e12); + int _e12; bar.InterlockedAnd(64, 5, _e12); tmp = _e12; - int _e15; bar.InterlockedXor(64, 5, _e15); + int _e15; bar.InterlockedOr(64, 5, _e15); tmp = _e15; - int _e18; bar.InterlockedMin(64, 5, _e18); + int _e18; bar.InterlockedXor(64, 5, _e18); tmp = _e18; - int _e21; bar.InterlockedMax(64, 5, _e21); + int _e21; bar.InterlockedMin(64, 5, _e21); tmp = _e21; - int _e24; bar.InterlockedExchange(64, 5, _e24); + int _e24; bar.InterlockedMax(64, 5, _e24); tmp = _e24; + int _e27; bar.InterlockedExchange(64, 5, _e27); + tmp = _e27; bar.Store(64, asuint(value)); return; } diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index fa0be834bf..cb9f88101a 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -58,18 +58,20 @@ kernel void atomics( int value = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed); int _e6 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e6; - int _e9 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e9 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e9; - int _e12 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e12 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e12; - int _e15 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e15 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e15; - int _e18 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e18 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e18; - int _e21 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e21 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e21; - int _e24 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); + int _e24 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed); tmp = _e24; + int _e27 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed); + tmp = _e27; metal::atomic_store_explicit(&bar.atom, value, metal::memory_order_relaxed); return; } diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index df869e495b..40698045e8 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 105 +; Bound: 107 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" @@ -139,24 +139,27 @@ OpBranch %85 %90 = OpAtomicIAdd %4 %91 %11 %89 %16 OpStore %82 %90 %93 = OpAccessChain %86 %27 %15 -%92 = OpAtomicAnd %4 %93 %11 %89 %16 +%92 = OpAtomicISub %4 %93 %11 %89 %16 OpStore %82 %92 %95 = OpAccessChain %86 %27 %15 -%94 = OpAtomicOr %4 %95 %11 %89 %16 +%94 = OpAtomicAnd %4 %95 %11 %89 %16 OpStore %82 %94 %97 = OpAccessChain %86 %27 %15 -%96 = OpAtomicXor %4 %97 %11 %89 %16 +%96 = OpAtomicOr %4 %97 %11 %89 %16 OpStore %82 %96 %99 = OpAccessChain %86 %27 %15 -%98 = OpAtomicSMin %4 %99 %11 %89 %16 +%98 = OpAtomicXor %4 %99 %11 %89 %16 OpStore %82 %98 %101 = OpAccessChain %86 %27 %15 -%100 = OpAtomicSMax %4 %101 %11 %89 %16 +%100 = OpAtomicSMin %4 %101 %11 %89 %16 OpStore %82 %100 %103 = OpAccessChain %86 %27 %15 -%102 = OpAtomicExchange %4 %103 %11 %89 %16 +%102 = OpAtomicSMax %4 %103 %11 %89 %16 OpStore %82 %102 -%104 = OpAccessChain %86 %27 %15 -OpAtomicStore %104 %11 %89 %88 +%105 = OpAccessChain %86 %27 %15 +%104 = OpAtomicExchange %4 %105 %11 %89 %16 +OpStore %82 %104 +%106 = OpAccessChain %86 %27 %15 +OpAtomicStore %106 %11 %89 %88 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index d9a00535a1..5eca81ff08 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -37,18 +37,20 @@ fn atomics() { let value: i32 = atomicLoad(&bar.atom); let _e6: i32 = atomicAdd(&bar.atom, 5); tmp = _e6; - let _e9: i32 = atomicAnd(&bar.atom, 5); + let _e9: i32 = atomicSub(&bar.atom, 5); tmp = _e9; - let _e12: i32 = atomicOr(&bar.atom, 5); + let _e12: i32 = atomicAnd(&bar.atom, 5); tmp = _e12; - let _e15: i32 = atomicXor(&bar.atom, 5); + let _e15: i32 = atomicOr(&bar.atom, 5); tmp = _e15; - let _e18: i32 = atomicMin(&bar.atom, 5); + let _e18: i32 = atomicXor(&bar.atom, 5); tmp = _e18; - let _e21: i32 = atomicMax(&bar.atom, 5); + let _e21: i32 = atomicMin(&bar.atom, 5); tmp = _e21; - let _e24: i32 = atomicExchange(&bar.atom, 5); + let _e24: i32 = atomicMax(&bar.atom, 5); tmp = _e24; + let _e27: i32 = atomicExchange(&bar.atom, 5); + tmp = _e27; atomicStore(&bar.atom, value); return; }