From f99671b830811d33dc1d1e234c4b66300ed007a7 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Thu, 23 Jun 2022 13:14:54 +0000 Subject: [PATCH] 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 Commit-Queue: Antonio Maiorano Kokoro: Kokoro --- src/tint/writer/hlsl/generator_impl.cc | 4 +- src/tint/writer/spirv/builder.cc | 3 +- .../1bd40a.wgsl.expected.spvasm | 26 ++++++------- .../63d8e6.wgsl.expected.spvasm | 26 ++++++------- .../83580d.wgsl.expected.hlsl | 39 +++++++++++++------ .../83580d.wgsl.expected.spvasm | 32 +++++++-------- .../e88938.wgsl.expected.hlsl | 39 +++++++++++++------ .../e88938.wgsl.expected.spvasm | 32 +++++++-------- 8 files changed, 110 insertions(+), 91 deletions(-) diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 3dfb74f0b0..de7a0beec1 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -1779,8 +1779,8 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, { // T 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 4368e3c006..c71f0fcb3d 100644 --- a/src/tint/writer/spirv/builder.cc +++ b/src/tint/writer/spirv/builder.cc @@ -3272,7 +3272,8 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call, 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 0c9313476f..0999969646 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 @@ SKIP: FAILED %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 62eac6119e..3723643514 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 @@ SKIP: FAILED %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 854c568161..a445a1d0ab 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 arg_0 : atomic; - -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 4f519e03e4..b163b6f12c 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 @@ SKIP: FAILED %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 ca63be97fb..c14dade47d 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 arg_0 : atomic; - -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 9bf039b6ee..b4e3102de6 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 @@ SKIP: FAILED %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 -