HLSL-IR: fix error: unary: no matching overload for 'operator - (u32)' (again)
My first fix handled this for polyfills:
https://dawn-review.googlesource.com/c/dawn/+/207475
This same fix is for DecomposeStorageAccess.
Bug: 368447426
Change-Id: I2e254a6c84ebe793659c5e5d86a25ee8d1b53141
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/210534
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: James Price <jrprice@google.com>
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
