diff --git a/src/tint/lang/hlsl/writer/builtin_test.cc b/src/tint/lang/hlsl/writer/builtin_test.cc
index fd0c82b..f6d8609 100644
--- a/src/tint/lang/hlsl/writer/builtin_test.cc
+++ b/src/tint/lang/hlsl/writer/builtin_test.cc
@@ -291,7 +291,7 @@
 RWByteAddressBuffer v : register(u0);
 void foo() {
   int v_1 = int(0);
-  v.InterlockedAdd(int(16u), -(int(123)), v_1);
+  v.InterlockedAdd(int(16u), (int(0) - int(123)), v_1);
   int x = v_1;
 }
 
@@ -314,7 +314,7 @@
 RWByteAddressBuffer v : register(u0);
 void foo() {
   int v_1 = int(0);
-  v.InterlockedAdd(int(0u), -(int(123)), v_1);
+  v.InterlockedAdd(int(0u), (int(0) - int(123)), v_1);
   int x = v_1;
 }
 
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc b/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
index 6505542..5366784 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
@@ -288,7 +288,7 @@
             auto* original_value = b.Var(ty.ptr(function, type));
             original_value->SetInitializer(b.Zero(type));
 
-            auto* val = b.Negation(type, args[1]);
+            auto* val = b.Subtract(type, b.Zero(type), args[1]);
             b.MemberCall<hlsl::ir::MemberBuiltinCall>(ty.void_(), BuiltinFn::kInterlockedAdd, var,
                                                       b.Convert(type, u32(offset)), val,
                                                       original_value);
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
index 036562b..6002ce0 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
@@ -1585,7 +1585,7 @@
 %foo = @fragment func():void {
   $B2: {
     %3:ptr<function, i32, read_write> = var, 0i
-    %4:i32 = negation 123i
+    %4:i32 = sub 0i, 123i
     %5:i32 = convert 16u
     %6:void = %v.InterlockedAdd %5, %4, %3
     %7:i32 = load %3
@@ -1633,7 +1633,7 @@
 %foo = @fragment func():void {
   $B2: {
     %3:ptr<function, i32, read_write> = var, 0i
-    %4:i32 = negation 123i
+    %4:i32 = sub 0i, 123i
     %5:i32 = convert 0u
     %6:void = %v.InterlockedAdd %5, %4, %3
     %7:i32 = load %3
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
index 36bc0fb..f35e34c 100644
--- a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
@@ -3,7 +3,7 @@
 void atomicSub_051100() {
   int res = int(0);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_9 = v;
   res = x_9;
 }
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
index 36bc0fb..f35e34c 100644
--- a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
@@ -3,7 +3,7 @@
 void atomicSub_051100() {
   int res = int(0);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_9 = v;
   res = x_9;
 }
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
index ac8d6e1..ee8fc2b 100644
--- a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
@@ -1,106 +1,27 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicSub_15bfc9() {
-  var res = 0u;
-  let x_9 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicSub_15bfc9() {
+  uint res = 0u;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_9 = v;
   res = x_9;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicSub_15bfc9 = func():void {
-  $B2: {
-    %res:ptr<function, u32, read_write> = var, 0u
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %x_9:u32 = let %8
-    store %res, %x_9
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %11:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %13:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %15:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %17:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
index ac8d6e1..ee8fc2b 100644
--- a/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
@@ -1,106 +1,27 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicSub_15bfc9() {
-  var res = 0u;
-  let x_9 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicSub_15bfc9() {
+  uint res = 0u;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_9 = v;
   res = x_9;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicSub_15bfc9 = func():void {
-  $B2: {
-    %res:ptr<function, u32, read_write> = var, 0u
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %x_9:u32 = let %8
-    store %res, %x_9
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %11:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %13:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %15:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %17:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
index 64c547c..88f10c7 100644
--- a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
@@ -3,7 +3,7 @@
 void atomicAdd_d32fe4() {
   int res = int(0);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_9 = v;
   res = x_9;
 }
diff --git a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
index 64c547c..88f10c7 100644
--- a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
@@ -3,7 +3,7 @@
 void atomicAdd_d32fe4() {
   int res = int(0);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_9 = v;
   res = x_9;
 }
diff --git a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
index a27f4ce..edfd362 100644
--- a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
@@ -1,106 +1,27 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicAdd_8a199a() {
-  var res = 0u;
-  let x_9 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicAdd_8a199a() {
+  uint res = 0u;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_9 = v;
   res = x_9;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicAdd_8a199a = func():void {
-  $B2: {
-    %res:ptr<function, u32, read_write> = var, 0u
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %x_9:u32 = let %8
-    store %res, %x_9
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %11:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %13:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %15:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %17:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
index a27f4ce..edfd362 100644
--- a/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/literal/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
@@ -1,106 +1,27 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicAdd_8a199a() {
-  var res = 0u;
-  let x_9 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicAdd_8a199a() {
+  uint res = 0u;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_9 = v;
   res = x_9;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicAdd_8a199a = func():void {
-  $B2: {
-    %res:ptr<function, u32, read_write> = var, 0u
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %x_9:u32 = let %8
-    store %res, %x_9
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %11:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %13:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %15:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %17:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
index 59f14e7..27e0297 100644
--- a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.dxc.hlsl
@@ -6,7 +6,7 @@
   arg_1 = int(1);
   int x_20 = arg_1;
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(x_20), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - x_20), v);
   int x_13 = v;
   res = x_13;
 }
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
index 59f14e7..27e0297 100644
--- a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_i32.spvasm.expected.ir.fxc.hlsl
@@ -6,7 +6,7 @@
   arg_1 = int(1);
   int x_20 = arg_1;
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(x_20), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - x_20), v);
   int x_13 = v;
   res = x_13;
 }
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
index 97f5aa6..62c0131 100644
--- a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.dxc.hlsl
@@ -1,113 +1,30 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicSub_15bfc9() {
-  var arg_1 = 0u;
-  var res = 0u;
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicSub_15bfc9() {
+  uint arg_1 = 0u;
+  uint res = 0u;
   arg_1 = 1u;
-  let x_18 = arg_1;
-  let x_13 = atomicSub(&(sb_rw.arg_0), x_18);
+  uint x_18 = arg_1;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - x_18), v);
+  uint x_13 = v;
   res = x_13;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :17:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %8:u32 = negation %x_18
-    ^^^^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicSub_15bfc9 = func():void {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 0u
-    %res:ptr<function, u32, read_write> = var, 0u
-    store %arg_1, 1u
-    %5:u32 = load %arg_1
-    %x_18:u32 = let %5
-    %7:ptr<function, u32, read_write> = var, 0u
-    %8:u32 = negation %x_18
-    %9:u32 = convert 0u
-    %10:void = %sb_rw.InterlockedAdd %9, %8, %7
-    %11:u32 = load %7
-    %x_13:u32 = let %11
-    store %res, %x_13
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %14:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %16:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %18:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %20:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
index 97f5aa6..62c0131 100644
--- a/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/atomicSub/storage_u32.spvasm.expected.ir.fxc.hlsl
@@ -1,113 +1,30 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicSub_15bfc9() {
-  var arg_1 = 0u;
-  var res = 0u;
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicSub_15bfc9() {
+  uint arg_1 = 0u;
+  uint res = 0u;
   arg_1 = 1u;
-  let x_18 = arg_1;
-  let x_13 = atomicSub(&(sb_rw.arg_0), x_18);
+  uint x_18 = arg_1;
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - x_18), v);
+  uint x_13 = v;
   res = x_13;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicSub_15bfc9();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :17:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %8:u32 = negation %x_18
-    ^^^^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicSub_15bfc9 = func():void {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 0u
-    %res:ptr<function, u32, read_write> = var, 0u
-    store %arg_1, 1u
-    %5:u32 = load %arg_1
-    %x_18:u32 = let %5
-    %7:ptr<function, u32, read_write> = var, 0u
-    %8:u32 = negation %x_18
-    %9:u32 = convert 0u
-    %10:void = %sb_rw.InterlockedAdd %9, %8, %7
-    %11:u32 = load %7
-    %x_13:u32 = let %11
-    store %res, %x_13
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %14:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %16:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %18:void = call %atomicSub_15bfc9
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %20:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
index b0af903..453c193 100644
--- a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.dxc.hlsl
@@ -5,7 +5,7 @@
   int res = int(0);
   arg_1 = int(1);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_13 = v;
   res = x_13;
 }
diff --git a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
index b0af903..453c193 100644
--- a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_i32.spvasm.expected.ir.fxc.hlsl
@@ -5,7 +5,7 @@
   int res = int(0);
   arg_1 = int(1);
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int x_13 = v;
   res = x_13;
 }
diff --git a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
index dc74ae9..db614b7 100644
--- a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.dxc.hlsl
@@ -1,110 +1,29 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicAdd_8a199a() {
-  var arg_1 = 0u;
-  var res = 0u;
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicAdd_8a199a() {
+  uint arg_1 = 0u;
+  uint res = 0u;
   arg_1 = 1u;
-  let x_13 = atomicSub(&(sb_rw.arg_0), 1u);
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_13 = v;
   res = x_13;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :15:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %6:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicAdd_8a199a = func():void {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 0u
-    %res:ptr<function, u32, read_write> = var, 0u
-    store %arg_1, 1u
-    %5:ptr<function, u32, read_write> = var, 0u
-    %6:u32 = negation 1u
-    %7:u32 = convert 0u
-    %8:void = %sb_rw.InterlockedAdd %7, %6, %5
-    %9:u32 = load %5
-    %x_13:u32 = let %9
-    store %res, %x_13
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %12:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %14:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %16:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %18:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
index dc74ae9..db614b7 100644
--- a/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/atomics/from_gen/var/spvAtomicDecrement/storage_u32.spvasm.expected.ir.fxc.hlsl
@@ -1,110 +1,29 @@
-SKIP: FAILED
 
-
-struct SB_RW_atomic {
-  /* @offset(0) */
-  arg_0 : atomic<u32>,
-}
-
-struct SB_RW {
-  /* @offset(0) */
-  arg_0 : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
-
-fn atomicAdd_8a199a() {
-  var arg_1 = 0u;
-  var res = 0u;
+RWByteAddressBuffer sb_rw : register(u0);
+void atomicAdd_8a199a() {
+  uint arg_1 = 0u;
+  uint res = 0u;
   arg_1 = 1u;
-  let x_13 = atomicSub(&(sb_rw.arg_0), 1u);
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint x_13 = v;
   res = x_13;
-  return;
 }
 
-fn fragment_main_1() {
+void fragment_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@fragment
-fn fragment_main() {
+void fragment_main() {
   fragment_main_1();
 }
 
-fn compute_main_1() {
+void compute_main_1() {
   atomicAdd_8a199a();
-  return;
 }
 
-@compute @workgroup_size(1i, 1i, 1i)
-fn compute_main() {
+[numthreads(1, 1, 1)]
+void compute_main() {
   compute_main_1();
 }
 
-Failed to generate: :15:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %6:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:10:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW_atomic = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-}
-
-%atomicAdd_8a199a = func():void {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 0u
-    %res:ptr<function, u32, read_write> = var, 0u
-    store %arg_1, 1u
-    %5:ptr<function, u32, read_write> = var, 0u
-    %6:u32 = negation 1u
-    %7:u32 = convert 0u
-    %8:void = %sb_rw.InterlockedAdd %7, %6, %5
-    %9:u32 = load %5
-    %x_13:u32 = let %9
-    store %res, %x_13
-    ret
-  }
-}
-%fragment_main_1 = func():void {
-  $B3: {
-    %12:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%fragment_main = @fragment func():void {
-  $B4: {
-    %14:void = call %fragment_main_1
-    ret
-  }
-}
-%compute_main_1 = func():void {
-  $B5: {
-    %16:void = call %atomicAdd_8a199a
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B6: {
-    %18:void = call %compute_main_1
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
index e8df965..fa8ba4f 100644
--- a/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
@@ -3,7 +3,7 @@
 RWByteAddressBuffer sb_rw : register(u1);
 int atomicSub_051100() {
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int res = v;
   return res;
 }
diff --git a/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
index e8df965..fa8ba4f 100644
--- a/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
@@ -3,7 +3,7 @@
 RWByteAddressBuffer sb_rw : register(u1);
 int atomicSub_051100() {
   int v = int(0);
-  sb_rw.InterlockedAdd(int(0u), -(int(1)), v);
+  sb_rw.InterlockedAdd(int(0u), (int(0) - int(1)), v);
   int res = v;
   return res;
 }
diff --git a/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
index 4e24ca0..b247b73 100644
--- a/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
@@ -1,82 +1,19 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-struct SB_RW {
-  arg_0 : atomic<u32>,
-}
-
-@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicSub_15bfc9() -> u32 {
-  var res : u32 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+RWByteAddressBuffer sb_rw : register(u1);
+uint atomicSub_15bfc9() {
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint res = v;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = atomicSub_15bfc9();
+void fragment_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = atomicSub_15bfc9();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:11:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 1)
-}
-
-%atomicSub_15bfc9 = func():u32 {
-  $B2: {
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %res:ptr<function, u32, read_write> = var, %8
-    %10:u32 = load %res
-    ret %10
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %12:u32 = call %atomicSub_15bfc9
-    %13:u32 = bitcast %12
-    %14:void = %prevent_dce.Store 0u, %13
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %16:u32 = call %atomicSub_15bfc9
-    %17:u32 = bitcast %16
-    %18:void = %prevent_dce.Store 0u, %17
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
index 4e24ca0..b247b73 100644
--- a/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/literal/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
@@ -1,82 +1,19 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-struct SB_RW {
-  arg_0 : atomic<u32>,
-}
-
-@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicSub_15bfc9() -> u32 {
-  var res : u32 = atomicSub(&(sb_rw.arg_0), 1u);
+RWByteAddressBuffer prevent_dce : register(u0);
+RWByteAddressBuffer sb_rw : register(u1);
+uint atomicSub_15bfc9() {
+  uint v = 0u;
+  sb_rw.InterlockedAdd(uint(0u), (0u - 1u), v);
+  uint res = v;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = atomicSub_15bfc9();
+void fragment_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = atomicSub_15bfc9();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-Failed to generate: :13:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %5:u32 = negation 1u
-    ^^^^^^^^^^^^^^^^^^^^
-
-:11:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 1)
-}
-
-%atomicSub_15bfc9 = func():u32 {
-  $B2: {
-    %4:ptr<function, u32, read_write> = var, 0u
-    %5:u32 = negation 1u
-    %6:u32 = convert 0u
-    %7:void = %sb_rw.InterlockedAdd %6, %5, %4
-    %8:u32 = load %4
-    %res:ptr<function, u32, read_write> = var, %8
-    %10:u32 = load %res
-    ret %10
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %12:u32 = call %atomicSub_15bfc9
-    %13:u32 = bitcast %12
-    %14:void = %prevent_dce.Store 0u, %13
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %16:u32 = call %atomicSub_15bfc9
-    %17:u32 = bitcast %16
-    %18:void = %prevent_dce.Store 0u, %17
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
index 4582684..0a866b0 100644
--- a/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.dxc.hlsl
@@ -4,7 +4,7 @@
 int atomicSub_051100() {
   int arg_1 = int(1);
   int v = int(0);
-  int v_1 = -(arg_1);
+  int v_1 = (int(0) - arg_1);
   sb_rw.InterlockedAdd(int(0u), v_1, v);
   int res = v;
   return res;
diff --git a/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
index 4582684..0a866b0 100644
--- a/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/atomicSub/051100.wgsl.expected.ir.fxc.hlsl
@@ -4,7 +4,7 @@
 int atomicSub_051100() {
   int arg_1 = int(1);
   int v = int(0);
-  int v_1 = -(arg_1);
+  int v_1 = (int(0) - arg_1);
   sb_rw.InterlockedAdd(int(0u), v_1, v);
   int res = v;
   return res;
diff --git a/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl b/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
index e0a00ab..03af56d 100644
--- a/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.dxc.hlsl
@@ -1,85 +1,21 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-struct SB_RW {
-  arg_0 : atomic<u32>,
-}
-
-@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicSub_15bfc9() -> u32 {
-  var arg_1 = 1u;
-  var res : u32 = atomicSub(&(sb_rw.arg_0), arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+RWByteAddressBuffer sb_rw : register(u1);
+uint atomicSub_15bfc9() {
+  uint arg_1 = 1u;
+  uint v = 0u;
+  uint v_1 = (0u - arg_1);
+  sb_rw.InterlockedAdd(uint(0u), v_1, v);
+  uint res = v;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = atomicSub_15bfc9();
+void fragment_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = atomicSub_15bfc9();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-Failed to generate: :15:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %7:u32 = negation %5
-    ^^^^^^^^^^^^^^^^^^^^
-
-:11:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 1)
-}
-
-%atomicSub_15bfc9 = func():u32 {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:u32 = load %arg_1
-    %6:ptr<function, u32, read_write> = var, 0u
-    %7:u32 = negation %5
-    %8:u32 = convert 0u
-    %9:void = %sb_rw.InterlockedAdd %8, %7, %6
-    %10:u32 = load %6
-    %res:ptr<function, u32, read_write> = var, %10
-    %12:u32 = load %res
-    ret %12
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %14:u32 = call %atomicSub_15bfc9
-    %15:u32 = bitcast %14
-    %16:void = %prevent_dce.Store 0u, %15
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %18:u32 = call %atomicSub_15bfc9
-    %19:u32 = bitcast %18
-    %20:void = %prevent_dce.Store 0u, %19
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1
diff --git a/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl b/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
index e0a00ab..03af56d 100644
--- a/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/builtins/gen/var/atomicSub/15bfc9.wgsl.expected.ir.fxc.hlsl
@@ -1,85 +1,21 @@
-SKIP: FAILED
 
-
-@group(0) @binding(0) var<storage, read_write> prevent_dce : u32;
-
-struct SB_RW {
-  arg_0 : atomic<u32>,
-}
-
-@group(0) @binding(1) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicSub_15bfc9() -> u32 {
-  var arg_1 = 1u;
-  var res : u32 = atomicSub(&(sb_rw.arg_0), arg_1);
+RWByteAddressBuffer prevent_dce : register(u0);
+RWByteAddressBuffer sb_rw : register(u1);
+uint atomicSub_15bfc9() {
+  uint arg_1 = 1u;
+  uint v = 0u;
+  uint v_1 = (0u - arg_1);
+  sb_rw.InterlockedAdd(uint(0u), v_1, v);
+  uint res = v;
   return res;
 }
 
-@fragment
-fn fragment_main() {
-  prevent_dce = atomicSub_15bfc9();
+void fragment_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-@compute @workgroup_size(1)
-fn compute_main() {
-  prevent_dce = atomicSub_15bfc9();
+[numthreads(1, 1, 1)]
+void compute_main() {
+  prevent_dce.Store(0u, atomicSub_15bfc9());
 }
 
-Failed to generate: :15:5 error: unary: no matching overload for 'operator - (u32)'
-
-2 candidate operators:
- • 'operator - (T  ✗ ) -> T' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
- • 'operator - (vecN<T>  ✗ ) -> vecN<T>' where:
-      ✗  'T' is 'f32', 'i32' or 'f16'
-
-    %7:u32 = negation %5
-    ^^^^^^^^^^^^^^^^^^^^
-
-:11:3 note: in block
-  $B2: {
-  ^^^
-
-note: # Disassembly
-SB_RW = struct @align(4) {
-  arg_0:atomic<u32> @offset(0)
-}
-
-$B1: {  # root
-  %prevent_dce:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
-  %sb_rw:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 1)
-}
-
-%atomicSub_15bfc9 = func():u32 {
-  $B2: {
-    %arg_1:ptr<function, u32, read_write> = var, 1u
-    %5:u32 = load %arg_1
-    %6:ptr<function, u32, read_write> = var, 0u
-    %7:u32 = negation %5
-    %8:u32 = convert 0u
-    %9:void = %sb_rw.InterlockedAdd %8, %7, %6
-    %10:u32 = load %6
-    %res:ptr<function, u32, read_write> = var, %10
-    %12:u32 = load %res
-    ret %12
-  }
-}
-%fragment_main = @fragment func():void {
-  $B3: {
-    %14:u32 = call %atomicSub_15bfc9
-    %15:u32 = bitcast %14
-    %16:void = %prevent_dce.Store 0u, %15
-    ret
-  }
-}
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B4: {
-    %18:u32 = call %atomicSub_15bfc9
-    %19:u32 = bitcast %18
-    %20:void = %prevent_dce.Store 0u, %19
-    ret
-  }
-}
-
-
-tint executable returned error: exit status 1