Implement atomicSub intrinsic

Polyfill this for HLSL using an atomic add with the operand negated.

Fixed: tint:1130
Change-Id: Ifa32d58973f1b48593ec0f6320f47f4358a5a3a9
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/62760
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/intrinsics.def

@@ -537,6 +537,7 @@ fn textureLoad(texture: texture_external, coords: vec2<i32>) -> vec4<f32>
 [[stage("fragment", "compute")]] fn atomicLoad<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>) -> T
 [[stage("fragment", "compute")]] fn atomicStore<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T)
 [[stage("fragment", "compute")]] fn atomicAdd<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
+[[stage("fragment", "compute")]] fn atomicSub<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
 [[stage("fragment", "compute")]] fn atomicMax<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
 [[stage("fragment", "compute")]] fn atomicMin<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
 [[stage("fragment", "compute")]] fn atomicAnd<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T

src/sem/intrinsic.cc

@@ -91,6 +91,7 @@ bool IsAtomicIntrinsic(IntrinsicType i) {
   return i == sem::IntrinsicType::kAtomicLoad ||
          i == sem::IntrinsicType::kAtomicStore ||
          i == sem::IntrinsicType::kAtomicAdd ||
+         i == sem::IntrinsicType::kAtomicSub ||
          i == sem::IntrinsicType::kAtomicMax ||
          i == sem::IntrinsicType::kAtomicMin ||
          i == sem::IntrinsicType::kAtomicAnd ||

src/sem/intrinsic_type.cc

@@ -303,6 +303,9 @@ IntrinsicType ParseIntrinsicType(const std::string& name) {
   if (name == "atomicAdd") {
     return IntrinsicType::kAtomicAdd;
   }
+  if (name == "atomicSub") {
+    return IntrinsicType::kAtomicSub;
+  }
   if (name == "atomicMax") {
     return IntrinsicType::kAtomicMax;
   }
@@ -513,6 +516,8 @@ const char* str(IntrinsicType i) {
       return "atomicStore";
     case IntrinsicType::kAtomicAdd:
       return "atomicAdd";
+    case IntrinsicType::kAtomicSub:
+      return "atomicSub";
     case IntrinsicType::kAtomicMax:
       return "atomicMax";
     case IntrinsicType::kAtomicMin:

src/sem/intrinsic_type.h

@@ -125,6 +125,7 @@ enum class IntrinsicType {
   kAtomicLoad,
   kAtomicStore,
   kAtomicAdd,
+  kAtomicSub,
   kAtomicMax,
   kAtomicMin,
   kAtomicAnd,

src/transform/decompose_memory_access.cc

@@ -238,6 +238,9 @@ DecomposeMemoryAccess::Intrinsic* IntrinsicAtomicFor(ProgramBuilder* builder,
     case sem::IntrinsicType::kAtomicAdd:
       op = DecomposeMemoryAccess::Intrinsic::Op::kAtomicAdd;
       break;
+    case sem::IntrinsicType::kAtomicSub:
+      op = DecomposeMemoryAccess::Intrinsic::Op::kAtomicSub;
+      break;
     case sem::IntrinsicType::kAtomicMax:
       op = DecomposeMemoryAccess::Intrinsic::Op::kAtomicMax;
       break;
@@ -723,6 +726,9 @@ std::string DecomposeMemoryAccess::Intrinsic::InternalName() const {
     case Op::kAtomicAdd:
       ss << "intrinsic_atomic_add_";
       break;
+    case Op::kAtomicSub:
+      ss << "intrinsic_atomic_sub_";
+      break;
     case Op::kAtomicMax:
       ss << "intrinsic_atomic_max_";
       break;

src/transform/decompose_memory_access.h

@@ -46,6 +46,7 @@ class DecomposeMemoryAccess
       kAtomicLoad,
       kAtomicStore,
       kAtomicAdd,
+      kAtomicSub,
       kAtomicMax,
       kAtomicMin,
       kAtomicAnd,

src/transform/decompose_memory_access_test.cc

@@ -1237,6 +1237,7 @@ fn main() {
   atomicStore(&sb.a, 123);
   ignore(atomicLoad(&sb.a));
   ignore(atomicAdd(&sb.a, 123));
+  ignore(atomicSub(&sb.a, 123));
   ignore(atomicMax(&sb.a, 123));
   ignore(atomicMin(&sb.a, 123));
   ignore(atomicAnd(&sb.a, 123));
@@ -1248,6 +1249,7 @@ fn main() {
   atomicStore(&sb.b, 123u);
   ignore(atomicLoad(&sb.b));
   ignore(atomicAdd(&sb.b, 123u));
+  ignore(atomicSub(&sb.b, 123u));
   ignore(atomicMax(&sb.b, 123u));
   ignore(atomicMin(&sb.b, 123u));
   ignore(atomicAnd(&sb.b, 123u));
@@ -1277,56 +1279,62 @@ fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_pa
 [[internal(intrinsic_atomic_add_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_max_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_sub_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_min_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_max_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_and_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_min_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_or_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_and_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_xor_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_or_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
-[[internal(intrinsic_atomic_exchange_storage_i32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_xor_storage_i32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
 
+[[internal(intrinsic_atomic_exchange_storage_i32), internal(disable_validation__function_has_no_body)]]
+fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
+
 [[internal(intrinsic_atomic_compare_exchange_weak_storage_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
+fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
 
 [[internal(intrinsic_atomic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32)
+fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32)
 
 [[internal(intrinsic_atomic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
+fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
 
 [[internal(intrinsic_atomic_add_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
-
-[[internal(intrinsic_atomic_max_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
-[[internal(intrinsic_atomic_min_storage_u32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_sub_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
-[[internal(intrinsic_atomic_and_storage_u32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_max_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
-[[internal(intrinsic_atomic_or_storage_u32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_min_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
-[[internal(intrinsic_atomic_xor_storage_u32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_and_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
-[[internal(intrinsic_atomic_exchange_storage_u32), internal(disable_validation__function_has_no_body)]]
+[[internal(intrinsic_atomic_or_storage_u32), internal(disable_validation__function_has_no_body)]]
 fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
 
+[[internal(intrinsic_atomic_xor_storage_u32), internal(disable_validation__function_has_no_body)]]
+fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
+
+[[internal(intrinsic_atomic_exchange_storage_u32), internal(disable_validation__function_has_no_body)]]
+fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
+
 [[internal(intrinsic_atomic_compare_exchange_weak_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
@@ -1339,17 +1347,19 @@ fn main() {
   ignore(tint_symbol_6(sb, 16u, 123));
   ignore(tint_symbol_7(sb, 16u, 123));
   ignore(tint_symbol_8(sb, 16u, 123));
-  ignore(tint_symbol_9(sb, 16u, 123, 345));
+  ignore(tint_symbol_9(sb, 16u, 123));
-  tint_symbol_10(sb, 20u, 123u);
+  ignore(tint_symbol_10(sb, 16u, 123, 345));
-  ignore(tint_symbol_11(sb, 20u));
+  tint_symbol_11(sb, 20u, 123u);
-  ignore(tint_symbol_12(sb, 20u, 123u));
+  ignore(tint_symbol_12(sb, 20u));
   ignore(tint_symbol_13(sb, 20u, 123u));
   ignore(tint_symbol_14(sb, 20u, 123u));
   ignore(tint_symbol_15(sb, 20u, 123u));
   ignore(tint_symbol_16(sb, 20u, 123u));
   ignore(tint_symbol_17(sb, 20u, 123u));
   ignore(tint_symbol_18(sb, 20u, 123u));
-  ignore(tint_symbol_19(sb, 20u, 123u, 345u));
+  ignore(tint_symbol_19(sb, 20u, 123u));
+  ignore(tint_symbol_20(sb, 20u, 123u));
+  ignore(tint_symbol_21(sb, 20u, 123u, 345u));
 }
 )";
 
@@ -1373,6 +1383,7 @@ fn main() {
   atomicStore(&(w.a), 123);
   ignore(atomicLoad(&(w.a)));
   ignore(atomicAdd(&(w.a), 123));
+  ignore(atomicSub(&(w.a), 123));
   ignore(atomicMax(&(w.a), 123));
   ignore(atomicMin(&(w.a), 123));
   ignore(atomicAnd(&(w.a), 123));
@@ -1383,6 +1394,7 @@ fn main() {
   atomicStore(&(w.b), 123u);
   ignore(atomicLoad(&(w.b)));
   ignore(atomicAdd(&(w.b), 123u));
+  ignore(atomicSub(&(w.b), 123u));
   ignore(atomicMax(&(w.b), 123u));
   ignore(atomicMin(&(w.b), 123u));
   ignore(atomicAnd(&(w.b), 123u));

src/writer/hlsl/generator_impl.cc

@@ -869,6 +869,7 @@ bool GeneratorImpl::EmitStorageBufferAccess(
     case Op::kAtomicLoad:
     case Op::kAtomicStore:
     case Op::kAtomicAdd:
+    case Op::kAtomicSub:
     case Op::kAtomicMax:
     case Op::kAtomicMin:
     case Op::kAtomicAnd:
@@ -930,7 +931,14 @@ bool GeneratorImpl::EmitStorageAtomicCall(
         }
         l << " = 0;";
       }
-      line(&buf) << "buffer." << hlsl << "(offset, value, original_value);";
+      {
+        auto l = line(&buf);
+        l << "buffer." << hlsl << "(offset, ";
+        if (intrinsic->op == Op::kAtomicSub) {
+          l << "-";
+        }
+        l << "value, original_value);";
+      }
       line(&buf) << "return original_value;";
       return name;
     };
@@ -939,6 +947,10 @@ bool GeneratorImpl::EmitStorageAtomicCall(
       case Op::kAtomicAdd:
         return rmw("atomicAdd", "InterlockedAdd");
 
+      case Op::kAtomicSub:
+        // Use add with the operand negated.
+        return rmw("atomicSub", "InterlockedAdd");
+
       case Op::kAtomicMax:
         return rmw("atomicMax", "InterlockedMax");
 
@@ -1130,6 +1142,10 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
         if (i > 0) {
           pre << ", ";
         }
+        if (i == 1 && intrinsic->Type() == sem::IntrinsicType::kAtomicSub) {
+          // Sub uses InterlockedAdd with the operand negated.
+          pre << "-";
+        }
         if (!EmitExpression(pre, arg)) {
           return false;
         }
@@ -1240,6 +1256,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
     }
 
     case sem::IntrinsicType::kAtomicAdd:
+    case sem::IntrinsicType::kAtomicSub:
       return call("InterlockedAdd");
 
     case sem::IntrinsicType::kAtomicMax:

src/writer/msl/generator_impl.cc

@@ -634,6 +634,9 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
     case sem::IntrinsicType::kAtomicAdd:
       return call("atomic_fetch_add_explicit", true);
 
+    case sem::IntrinsicType::kAtomicSub:
+      return call("atomic_fetch_sub_explicit", true);
+
     case sem::IntrinsicType::kAtomicMax:
       return call("atomic_fetch_max_explicit", true);
 

src/writer/spirv/builder.cc

@@ -3151,6 +3151,15 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
                                                            semantics,
                                                            value,
                                                        });
+    case sem::IntrinsicType::kAtomicSub:
+      return push_function_inst(spv::Op::OpAtomicISub, {
+                                                           result_type,
+                                                           result_id,
+                                                           pointer,
+                                                           memory,
+                                                           semantics,
+                                                           value,
+                                                       });
     case sem::IntrinsicType::kAtomicMax:
       return push_function_inst(
           is_value_signed() ? spv::Op::OpAtomicSMax : spv::Op::OpAtomicUMax,

test/intrinsics/gen/atomicSub/051100.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[block]]
+struct SB_RW {
+  arg_0: atomic<i32>;
+};
+[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
+
+// fn atomicSub(ptr<storage, atomic<i32>, read_write>, i32) -> i32
+fn atomicSub_051100() {
+  var res: i32 = atomicSub(&sb_rw.arg_0, 1);
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  atomicSub_051100();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_051100();
+}

test/intrinsics/gen/atomicSub/051100.wgsl.expected.hlsl

@@ -0,0 +1,22 @@
+int atomicSub_1(RWByteAddressBuffer buffer, uint offset, int value) {
+  int original_value = 0;
+  buffer.InterlockedAdd(offset, -value, original_value);
+  return original_value;
+}
+
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+void atomicSub_051100() {
+  int res = atomicSub_1(sb_rw, 0u, 1);
+}
+
+void fragment_main() {
+  atomicSub_051100();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  atomicSub_051100();
+  return;
+}

test/intrinsics/gen/atomicSub/051100.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB_RW {
+  /* 0x0000 */ atomic_int arg_0;
+};
+
+void atomicSub_051100(device SB_RW& sb_rw) {
+  int res = atomic_fetch_sub_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
+}
+
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicSub_051100(sb_rw);
+  return;
+}
+
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicSub_051100(sb_rw);
+  return;
+}
+

test/intrinsics/gen/atomicSub/051100.wgsl.expected.spvasm

@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW "SB_RW"
+               OpMemberName %SB_RW 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicSub_051100 "atomicSub_051100"
+               OpName %res "res"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW Block
+               OpMemberDecorate %SB_RW 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+        %int = OpTypeInt 32 1
+      %SB_RW = OpTypeStruct %int
+%_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %19 = OpConstantNull %int
+%atomicSub_051100 = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %19
+         %15 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+          %9 = OpAtomicISub %int %15 %uint_1 %uint_0 %int_1
+               OpStore %res %9
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %atomicSub_051100
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %atomicSub_051100
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/atomicSub/051100.wgsl.expected.wgsl

@@ -0,0 +1,20 @@
+[[block]]
+struct SB_RW {
+  arg_0 : atomic<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
+
+fn atomicSub_051100() {
+  var res : i32 = atomicSub(&(sb_rw.arg_0), 1);
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  atomicSub_051100();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_051100();
+}

test/intrinsics/gen/atomicSub/0d26c2.wgsl

@@ -0,0 +1,35 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: atomic<u32>;
+
+// fn atomicSub(ptr<workgroup, atomic<u32>, read_write>, u32) -> u32
+fn atomicSub_0d26c2() {
+  var res: u32 = atomicSub(&arg_0, 1u);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_0d26c2();
+}

test/intrinsics/gen/atomicSub/0d26c2.wgsl.expected.hlsl

@@ -0,0 +1,26 @@
+groupshared uint arg_0;
+
+void atomicSub_0d26c2() {
+  uint atomic_result = 0u;
+  InterlockedAdd(arg_0, -1u, atomic_result);
+  uint res = atomic_result;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    uint atomic_result_1 = 0u;
+    InterlockedExchange(arg_0, 0u, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  atomicSub_0d26c2();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}

test/intrinsics/gen/atomicSub/0d26c2.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void atomicSub_0d26c2(threadgroup atomic_uint* const tint_symbol) {
+  uint res = atomic_fetch_sub_explicit(tint_symbol, 1u, memory_order_relaxed);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
+  {
+    atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicSub_0d26c2(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  compute_main_inner(local_invocation_index, &(tint_symbol_2));
+  return;
+}
+

test/intrinsics/gen/atomicSub/0d26c2.wgsl.expected.spvasm

@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicSub_0d26c2 "atomicSub_0d26c2"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+      %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+         %18 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicSub_0d26c2 = OpFunction %void None %6
+          %9 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %17
+         %10 = OpAtomicISub %uint %arg_0 %uint_2 %uint_0 %uint_1
+               OpStore %res %10
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %18
+%local_invocation_index = OpFunctionParameter %uint
+         %21 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %17
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %26 = OpFunctionCall %void %atomicSub_0d26c2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %6
+         %28 = OpLabel
+         %30 = OpLoad %uint %local_invocation_index_1
+         %29 = OpFunctionCall %void %compute_main_inner %30
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/atomicSub/0d26c2.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<u32>;
+
+fn atomicSub_0d26c2() {
+  var res : u32 = atomicSub(&(arg_0), 1u);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_0d26c2();
+}

test/intrinsics/gen/atomicSub/15bfc9.wgsl

@@ -0,0 +1,44 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+[[block]]
+struct SB_RW {
+  arg_0: atomic<u32>;
+};
+[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
+
+// fn atomicSub(ptr<storage, atomic<u32>, read_write>, u32) -> u32
+fn atomicSub_15bfc9() {
+  var res: u32 = atomicSub(&sb_rw.arg_0, 1u);
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  atomicSub_15bfc9();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_15bfc9();
+}

test/intrinsics/gen/atomicSub/15bfc9.wgsl.expected.hlsl

@@ -0,0 +1,22 @@
+uint atomicSub_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, -value, original_value);
+  return original_value;
+}
+
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+void atomicSub_15bfc9() {
+  uint res = atomicSub_1(sb_rw, 0u, 1u);
+}
+
+void fragment_main() {
+  atomicSub_15bfc9();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  atomicSub_15bfc9();
+  return;
+}

test/intrinsics/gen/atomicSub/15bfc9.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB_RW {
+  /* 0x0000 */ atomic_uint arg_0;
+};
+
+void atomicSub_15bfc9(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_sub_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
+}
+
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicSub_15bfc9(sb_rw);
+  return;
+}
+
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicSub_15bfc9(sb_rw);
+  return;
+}
+

test/intrinsics/gen/atomicSub/15bfc9.wgsl.expected.spvasm

@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW "SB_RW"
+               OpMemberName %SB_RW 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicSub_15bfc9 "atomicSub_15bfc9"
+               OpName %res "res"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW Block
+               OpMemberDecorate %SB_RW 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+       %uint = OpTypeInt 32 0
+      %SB_RW = OpTypeStruct %uint
+%_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+%atomicSub_15bfc9 = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %17
+         %14 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
+          %9 = OpAtomicISub %uint %14 %uint_1 %uint_0 %uint_1
+               OpStore %res %9
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %atomicSub_15bfc9
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %atomicSub_15bfc9
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/atomicSub/15bfc9.wgsl.expected.wgsl

@@ -0,0 +1,20 @@
+[[block]]
+struct SB_RW {
+  arg_0 : atomic<u32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
+
+fn atomicSub_15bfc9() {
+  var res : u32 = atomicSub(&(sb_rw.arg_0), 1u);
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  atomicSub_15bfc9();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_15bfc9();
+}

test/intrinsics/gen/atomicSub/77883a.wgsl

@@ -0,0 +1,35 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+var<workgroup> arg_0: atomic<i32>;
+
+// fn atomicSub(ptr<workgroup, atomic<i32>, read_write>, i32) -> i32
+fn atomicSub_77883a() {
+  var res: i32 = atomicSub(&arg_0, 1);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_77883a();
+}

test/intrinsics/gen/atomicSub/77883a.wgsl.expected.hlsl

@@ -0,0 +1,26 @@
+groupshared int arg_0;
+
+void atomicSub_77883a() {
+  int atomic_result = 0;
+  InterlockedAdd(arg_0, -1, atomic_result);
+  int res = atomic_result;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    int atomic_result_1 = 0;
+    InterlockedExchange(arg_0, 0, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  atomicSub_77883a();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}

test/intrinsics/gen/atomicSub/77883a.wgsl.expected.msl

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void atomicSub_77883a(threadgroup atomic_int* const tint_symbol) {
+  int res = atomic_fetch_sub_explicit(tint_symbol, 1, memory_order_relaxed);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+  {
+    atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicSub_77883a(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  compute_main_inner(local_invocation_index, &(tint_symbol_2));
+  return;
+}
+

test/intrinsics/gen/atomicSub/77883a.wgsl.expected.spvasm

@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicSub_77883a "atomicSub_77883a"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+      %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %18 = OpConstantNull %int
+         %19 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicSub_77883a = OpFunction %void None %7
+         %10 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %18
+         %11 = OpAtomicISub %int %arg_0 %uint_2 %uint_0 %int_1
+               OpStore %res %11
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %19
+%local_invocation_index = OpFunctionParameter %uint
+         %22 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %18
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %27 = OpFunctionCall %void %atomicSub_77883a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %7
+         %29 = OpLabel
+         %31 = OpLoad %uint %local_invocation_index_1
+         %30 = OpFunctionCall %void %compute_main_inner %31
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/atomicSub/77883a.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<i32>;
+
+fn atomicSub_77883a() {
+  var res : i32 = atomicSub(&(arg_0), 1);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  atomicSub_77883a();
+}