fix(glsl-out): add parens. to atomicSub negation

CHANGELOG.md

@@ -99,6 +99,7 @@ Bottom level categories:
 #### GLES
 
 - Fixed signed integer `%` (and `%=`) returning the wrong result for negative operands in the GLSL (OpenGL/GLES) backend, e.g. `-1 % 768` yielding `255` instead of `-1`. GLSL's `%` is undefined when either operand is negative, so signed remainder is now lowered as `a - b * (a / b)`, matching the SPIR-V, HLSL, and Metal backends. By @mstampfli in [#9687](https://github.com/gfx-rs/wgpu/pull/9687).
+- Fix negative argument for `atomicSub` yielding incorrect GLSL. By @ErichDonGubler in [#9924](https://github.com/gfx-rs/wgpu/pull/9924).
 
 ### Dependency Updates
 

naga/src/back/glsl/writer.rs

@@ -2128,9 +2128,15 @@ impl<'a, W: Write> Writer<'a, W> {
                         write!(self.out, ", ")?;
                         if let crate::AtomicFunction::Subtract = *fun {
                             // Emulate `atomicSub` with `atomicAdd` by negating the value.
-                            write!(self.out, "-")?;
+                            //
+                            // Make sure we use a surrounding parenthetical to avoid accidentally
+                            // producing a prefix increment (`--`).
+                            write!(self.out, "-(")?;
                         }
                         self.write_expr(value, ctx)?;
+                        if let crate::AtomicFunction::Subtract = *fun {
+                            write!(self.out, ")")?;
+                        }
                         writeln!(self.out, ");")?;
                     }
                 }

naga/tests/in/wgsl/9883-atomic_sub-negative.toml

@@ -0,0 +1 @@
+targets = "SPIRV | METAL | GLSL | HLSL | WGSL"

naga/tests/in/wgsl/9883-atomic_sub-negative.wgsl

@@ -0,0 +1,7 @@
+var<workgroup> a: atomic<i32>;
+
+@compute @workgroup_size(1)
+fn main() {
+    //  In GLSL, we accidentally printed `atomicAdd(a, --1)`. Oops!
+    let x = atomicSub(&a, -1i);
+}

naga/tests/out/glsl/wgsl-9883-atomic_sub-negative.main.Compute.glsl

@@ -0,0 +1,20 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+shared int a;
+
+
+void main() {
+    if (gl_LocalInvocationID == uvec3(0u)) {
+        a = 0;
+    }
+    memoryBarrierShared();
+    barrier();
+    int _e2 = atomicAdd(a, -(-1));
+    return;
+}
+

naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl

@@ -69,14 +69,14 @@ void main() {
     int _e80 = atomicAdd(workgroup_struct.atomic_arr[1], 1);
     memoryBarrierShared();
     barrier();
-    uint _e83 = atomicAdd(_group_0_binding_0_cs, -1u);
-    int _e87 = atomicAdd(_group_0_binding_1_cs[1], -1);
-    uint _e91 = atomicAdd(_group_0_binding_2_cs.atomic_scalar, -1u);
-    int _e96 = atomicAdd(_group_0_binding_2_cs.atomic_arr[1], -1);
-    uint _e99 = atomicAdd(workgroup_atomic_scalar, -1u);
-    int _e103 = atomicAdd(workgroup_atomic_arr[1], -1);
-    uint _e107 = atomicAdd(workgroup_struct.atomic_scalar, -1u);
-    int _e112 = atomicAdd(workgroup_struct.atomic_arr[1], -1);
+    uint _e83 = atomicAdd(_group_0_binding_0_cs, -(1u));
+    int _e87 = atomicAdd(_group_0_binding_1_cs[1], -(1));
+    uint _e91 = atomicAdd(_group_0_binding_2_cs.atomic_scalar, -(1u));
+    int _e96 = atomicAdd(_group_0_binding_2_cs.atomic_arr[1], -(1));
+    uint _e99 = atomicAdd(workgroup_atomic_scalar, -(1u));
+    int _e103 = atomicAdd(workgroup_atomic_arr[1], -(1));
+    uint _e107 = atomicAdd(workgroup_struct.atomic_scalar, -(1u));
+    int _e112 = atomicAdd(workgroup_struct.atomic_arr[1], -(1));
     memoryBarrierShared();
     barrier();
     uint _e115 = atomicMax(_group_0_binding_0_cs, 1u);

naga/tests/out/hlsl/wgsl-9883-atomic_sub-negative.hlsl

@@ -0,0 +1,12 @@
+groupshared int a;
+
+[numthreads(1, 1, 1)]
+void main(uint local_invocation_index : SV_GroupIndex)
+{
+    if (local_invocation_index == 0) {
+        a = (int)0;
+    }
+    GroupMemoryBarrierWithGroupSync();
+    int _e2; InterlockedAdd(a, -int(-1), _e2);
+    return;
+}

naga/tests/out/hlsl/wgsl-9883-atomic_sub-negative.ron

@@ -0,0 +1,16 @@
+(
+    vertex:[
+    ],
+    fragment:[
+    ],
+    compute:[
+        (
+            entry_point:"main",
+            target_profile:"cs_5_1",
+        ),
+    ],
+    task:[
+    ],
+    mesh:[
+    ],
+)

naga/tests/out/msl/wgsl-9883-atomic_sub-negative.metal

@@ -0,0 +1,18 @@
+// language: metal1.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+
+kernel void main_(
+  uint __local_invocation_index [[thread_index_in_threadgroup]]
+, threadgroup metal::atomic_int& a
+) {
+    if (__local_invocation_index == 0u) {
+        metal::atomic_store_explicit(&a, 0, metal::memory_order_relaxed);
+    }
+    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    int _e2 = metal::atomic_fetch_sub_explicit(&a, -1, metal::memory_order_relaxed);
+    return;
+}

naga/tests/out/spv/wgsl-9883-atomic_sub-negative.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 26
+OpCapability Shader
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %7 "main" %12
+OpExecutionMode %7 LocalSize 1 1 1
+OpDecorate %12 BuiltIn LocalInvocationIndex
+%2 = OpTypeVoid
+%3 = OpTypeInt 32 1
+%5 = OpTypePointer Workgroup %3
+%4 = OpVariable  %5  Workgroup
+%8 = OpTypeFunction %2
+%9 = OpConstant  %3  -1
+%11 = OpConstantNull  %3
+%13 = OpTypeInt 32 0
+%14 = OpTypePointer Input %13
+%12 = OpVariable  %14  Input
+%16 = OpConstant  %13  0
+%18 = OpTypeBool
+%21 = OpConstant  %13  2
+%22 = OpConstant  %13  264
+%25 = OpConstant  %3  2
+%7 = OpFunction  %2  None %8
+%6 = OpLabel
+OpBranch %10
+%10 = OpLabel
+%15 = OpLoad  %13  %12
+%17 = OpIEqual  %18  %15 %16
+OpSelectionMerge %19 None
+OpBranchConditional %17 %20 %19
+%20 = OpLabel
+OpStore %4 %11
+OpBranch %19
+%19 = OpLabel
+OpControlBarrier %21 %21 %22
+OpBranch %23
+%23 = OpLabel
+%24 = OpAtomicISub  %3  %4 %25 %16 %9
+OpReturn
+OpFunctionEnd

naga/tests/out/spv/wgsl-9883-atomic_sub-negative.spvasm.glsl

@@ -0,0 +1,15 @@
+#version 460
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+shared int _4;
+
+void main()
+{
+    if (gl_LocalInvocationIndex == 0u)
+    {
+        _4 = 0;
+    }
+    barrier();
+    int _24 = atomicAdd(_4, -(-1));
+}
+

naga/tests/out/wgsl/wgsl-9883-atomic_sub-negative.wgsl

@@ -0,0 +1,7 @@
+var<workgroup> a: atomic<i32>;
+
+@compute @workgroup_size(1, 1, 1) 
+fn main() {
+    let _e2 = atomicSub((&a), -1i);
+    return;
+}