tint: spir-v reader: fix atomicCompareExchangeWeak with var comparator
Also fix HLSL generator to unwrap the ref type when emitting the
comparator value.
Bug: tint:1185
Change-Id: I01d04ca6357e72fd5ead0f25012ab39794e65da5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94522
Reviewed-by: Ben Clayton <bclayton@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 3dfb74f..de7a0be 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1779,8 +1779,8 @@
{ // T compare_value = <compare_value>;
auto pre = line();
- if (!EmitTypeAndName(pre, TypeOf(compare_value), ast::StorageClass::kNone,
- ast::Access::kUndefined, compare)) {
+ if (!EmitTypeAndName(pre, TypeOf(compare_value)->UnwrapRef(),
+ ast::StorageClass::kNone, ast::Access::kUndefined, compare)) {
return false;
}
pre << " = ";
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 4368e3c..c71f0fc 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -3272,7 +3272,8 @@
value,
});
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
- auto comparator = GenerateExpression(call->Arguments()[1]->Declaration());
+ auto comparator =
+ GenerateExpressionWithLoadIfNeeded(call->Arguments()[1]->Declaration());
if (comparator == 0) {
return false;
}
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
index 0c93134..0999969 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 35
+; Bound: 36
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -45,32 +43,30 @@
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
- %28 = OpConstantNull %__atomic_compare_exchange_resulti32
+ %29 = OpConstantNull %__atomic_compare_exchange_resulti32
%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5
%8 = OpLabel
%arg_1 = OpVariable %_ptr_Function_int Function %12
%arg_2 = OpVariable %_ptr_Function_int Function %12
- %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
+ %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %29
OpStore %arg_1 %int_1
OpStore %arg_2 %int_1
%22 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
%23 = OpLoad %int %arg_2
- %24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
- %25 = OpIEqual %bool %24 %23
- %14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
+ %24 = OpLoad %int %arg_1
+ %25 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %24
+ %26 = OpIEqual %bool %25 %23
+ %14 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %25 %26
OpStore %res %14
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %5
- %30 = OpLabel
- %31 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %5
- %33 = OpLabel
- %34 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+ %34 = OpLabel
+ %35 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
OpReturn
OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
- %24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %23 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
index 62eac61..3723643 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 33
+; Bound: 34
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -43,32 +41,30 @@
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
- %26 = OpConstantNull %__atomic_compare_exchange_resultu32
+ %27 = OpConstantNull %__atomic_compare_exchange_resultu32
%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5
%8 = OpLabel
%arg_1 = OpVariable %_ptr_Function_uint Function %12
%arg_2 = OpVariable %_ptr_Function_uint Function %12
- %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
+ %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
OpStore %arg_1 %uint_1
OpStore %arg_2 %uint_1
%20 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
%21 = OpLoad %uint %arg_2
- %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
- %23 = OpIEqual %bool %22 %21
- %14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
+ %22 = OpLoad %uint %arg_1
+ %23 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %22
+ %24 = OpIEqual %bool %23 %21
+ %14 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
OpStore %res %14
OpReturn
OpFunctionEnd
%fragment_main = OpFunction %void None %5
- %28 = OpLabel
- %29 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+ %29 = OpLabel
+ %30 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %5
- %31 = OpLabel
- %32 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+ %32 = OpLabel
+ %33 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
OpReturn
OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
- %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
index 854c568..a445a1d 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
@@ -1,17 +1,34 @@
-SKIP: FAILED
+struct atomic_compare_exchange_resultu32 {
+ uint old_value;
+ bool exchanged;
+};
+groupshared uint arg_0;
-
-var<workgroup> arg_0 : atomic<u32>;
-
-fn atomicCompareExchangeWeak_83580d() {
- var arg_1 = 1u;
- var arg_2 = 1u;
- var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
+void atomicCompareExchangeWeak_83580d() {
+ uint arg_1 = 1u;
+ uint arg_2 = 1u;
+ atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0;
+ uint atomic_compare_value = arg_1;
+ InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
+ atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+ atomic_compare_exchange_resultu32 res = atomic_result;
}
-@compute @workgroup_size(1)
-fn compute_main() {
+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();
atomicCompareExchangeWeak_83580d();
}
-Failed to generate: error: unknown type in EmitType
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
index 4f519e0..b163b6f 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 40
+; Bound: 41
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -39,37 +37,35 @@
%uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
- %26 = OpConstantNull %__atomic_compare_exchange_resultu32
- %27 = OpTypeFunction %void %uint
+ %27 = OpConstantNull %__atomic_compare_exchange_resultu32
+ %28 = OpTypeFunction %void %uint
%uint_264 = OpConstant %uint 264
%atomicCompareExchangeWeak_83580d = OpFunction %void None %6
%9 = OpLabel
%arg_1 = OpVariable %_ptr_Function_uint Function %13
%arg_2 = OpVariable %_ptr_Function_uint Function %13
- %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %26
+ %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27
OpStore %arg_1 %uint_1
OpStore %arg_2 %uint_1
%21 = OpLoad %uint %arg_2
- %22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
- %23 = OpIEqual %bool %22 %21
- %15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23
+ %22 = OpLoad %uint %arg_1
+ %23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %22
+ %24 = OpIEqual %bool %23 %21
+ %15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24
OpStore %res %15
OpReturn
OpFunctionEnd
-%compute_main_inner = OpFunction %void None %27
+%compute_main_inner = OpFunction %void None %28
%local_invocation_index = OpFunctionParameter %uint
- %30 = OpLabel
+ %31 = OpLabel
OpAtomicStore %arg_0 %uint_2 %uint_0 %13
OpControlBarrier %uint_2 %uint_2 %uint_264
- %35 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
+ %36 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %6
- %37 = OpLabel
- %39 = OpLoad %uint %local_invocation_index_1
- %38 = OpFunctionCall %void %compute_main_inner %39
+ %38 = OpLabel
+ %40 = OpLoad %uint %local_invocation_index_1
+ %39 = OpFunctionCall %void %compute_main_inner %40
OpReturn
OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
- %22 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %arg_1
-
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
index ca63be9..c14dade 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
@@ -1,17 +1,34 @@
-SKIP: FAILED
+struct atomic_compare_exchange_resulti32 {
+ int old_value;
+ bool exchanged;
+};
+groupshared int arg_0;
-
-var<workgroup> arg_0 : atomic<i32>;
-
-fn atomicCompareExchangeWeak_e88938() {
- var arg_1 = 1;
- var arg_2 = 1;
- var res = atomicCompareExchangeWeak(&(arg_0), arg_1, arg_2);
+void atomicCompareExchangeWeak_e88938() {
+ int arg_1 = 1;
+ int arg_2 = 1;
+ atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0;
+ int atomic_compare_value = arg_1;
+ InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
+ atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+ atomic_compare_exchange_resulti32 res = atomic_result;
}
-@compute @workgroup_size(1)
-fn compute_main() {
+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();
atomicCompareExchangeWeak_e88938();
}
-Failed to generate: error: unknown type in EmitType
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
index 9bf039b..b4e3102 100644
--- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
@@ -1,9 +1,7 @@
-SKIP: FAILED
-
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 41
+; Bound: 42
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -40,37 +38,35 @@
%uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
- %27 = OpConstantNull %__atomic_compare_exchange_resulti32
- %28 = OpTypeFunction %void %uint
+ %28 = OpConstantNull %__atomic_compare_exchange_resulti32
+ %29 = OpTypeFunction %void %uint
%uint_264 = OpConstant %uint 264
%atomicCompareExchangeWeak_e88938 = OpFunction %void None %7
%10 = OpLabel
%arg_1 = OpVariable %_ptr_Function_int Function %14
%arg_2 = OpVariable %_ptr_Function_int Function %14
- %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %27
+ %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28
OpStore %arg_1 %int_1
OpStore %arg_2 %int_1
%22 = OpLoad %int %arg_2
- %23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
- %24 = OpIEqual %bool %23 %22
- %16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %23 %24
+ %23 = OpLoad %int %arg_1
+ %24 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %23
+ %25 = OpIEqual %bool %24 %22
+ %16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25
OpStore %res %16
OpReturn
OpFunctionEnd
-%compute_main_inner = OpFunction %void None %28
+%compute_main_inner = OpFunction %void None %29
%local_invocation_index = OpFunctionParameter %uint
- %31 = OpLabel
+ %32 = OpLabel
OpAtomicStore %arg_0 %uint_2 %uint_0 %14
OpControlBarrier %uint_2 %uint_2 %uint_264
- %36 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
+ %37 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
OpReturn
OpFunctionEnd
%compute_main = OpFunction %void None %7
- %38 = OpLabel
- %40 = OpLoad %uint %local_invocation_index_1
- %39 = OpFunctionCall %void %compute_main_inner %40
+ %39 = OpLabel
+ %41 = OpLoad %uint %local_invocation_index_1
+ %40 = OpFunctionCall %void %compute_main_inner %41
OpReturn
OpFunctionEnd
-1:1: AtomicCompareExchange: expected Comparator to be of type Result Type
- %23 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %arg_1
-
