diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc index eb8795f714..34d550b8eb 100644 --- a/src/tint/writer/spirv/builder.cc +++ b/src/tint/writer/spirv/builder.cc @@ -3161,7 +3161,7 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call, return false; } - auto* value_sem_type = TypeOf(call->Arguments()[2]->Declaration()); + auto* value_sem_type = call->Target()->Signature().parameters[2]->Type(); auto value_type = GenerateTypeIfNeeded(value_sem_type); if (value_type == 0) { diff --git a/test/tint/bug/tint/1573.wgsl b/test/tint/bug/tint/1573.wgsl new file mode 100644 index 0000000000..850e79c7b3 --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl @@ -0,0 +1,8 @@ +@group(0) @binding(0) +var a : atomic; + +@stage(compute) @workgroup_size(16) +fn main() { + var value = 42u; + let result = atomicCompareExchangeWeak(&a, 0u, value); +} diff --git a/test/tint/bug/tint/1573.wgsl.expected.glsl b/test/tint/bug/tint/1573.wgsl.expected.glsl new file mode 100644 index 0000000000..013839ef12 --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl.expected.glsl @@ -0,0 +1,28 @@ +#version 310 es + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; + + +struct a_block { + uint inner; +}; + +layout(binding = 0, std430) buffer a_block_1 { + uint inner; +} a; +void tint_symbol() { + uint value = 42u; + atomic_compare_exchange_resultu32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(a.inner, 0u, value); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u; + atomic_compare_exchange_resultu32 result = atomic_compare_result; +} + +layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1573.wgsl.expected.hlsl b/test/tint/bug/tint/1573.wgsl.expected.hlsl new file mode 100644 index 0000000000..a19a5806e1 --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl.expected.hlsl @@ -0,0 +1,21 @@ +RWByteAddressBuffer a : register(u0, space0); + +struct atomic_compare_exchange_weak_ret_type { + uint old_value; + bool exchanged; +}; + +atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) { + atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; + buffer.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +[numthreads(16, 1, 1)] +void main() { + uint value = 42u; + const atomic_compare_exchange_weak_ret_type result = tint_atomicCompareExchangeWeak(a, 0u, 0u, value); + return; +} diff --git a/test/tint/bug/tint/1573.wgsl.expected.msl b/test/tint/bug/tint/1573.wgsl.expected.msl new file mode 100644 index 0000000000..c284e1a5ea --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl.expected.msl @@ -0,0 +1,21 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; +template +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { + T old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +kernel void tint_symbol(device atomic_uint* tint_symbol_1 [[buffer(0)]]) { + uint value = 42u; + atomic_compare_exchange_resultu32 const result = atomicCompareExchangeWeak_1(tint_symbol_1, 0u, value); + return; +} + diff --git a/test/tint/bug/tint/1573.wgsl.expected.spvasm b/test/tint/bug/tint/1573.wgsl.expected.spvasm new file mode 100644 index 0000000000..1d2672bc88 --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl.expected.spvasm @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 24 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 16 1 1 + OpName %a_block "a_block" + OpMemberName %a_block 0 "inner" + OpName %a "a" + OpName %main "main" + OpName %value "value" + OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" + OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpDecorate %a_block Block + OpMemberDecorate %a_block 0 Offset 0 + OpDecorate %a DescriptorSet 0 + OpDecorate %a Binding 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + %uint = OpTypeInt 32 0 + %a_block = OpTypeStruct %uint +%_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block + %a = OpVariable %_ptr_StorageBuffer_a_block StorageBuffer + %void = OpTypeVoid + %5 = OpTypeFunction %void + %uint_42 = OpConstant %uint 42 +%_ptr_Function_uint = OpTypePointer Function %uint + %12 = OpConstantNull %uint + %bool = OpTypeBool +%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool + %uint_1 = OpConstant %uint 1 + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %main = OpFunction %void None %5 + %8 = OpLabel + %value = OpVariable %_ptr_Function_uint Function %12 + OpStore %value %uint_42 + %20 = OpAccessChain %_ptr_StorageBuffer_uint %a %uint_0 + %21 = OpLoad %uint %value + %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %uint_0 + %23 = OpIEqual %bool %22 %21 + %13 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1573.wgsl.expected.wgsl b/test/tint/bug/tint/1573.wgsl.expected.wgsl new file mode 100644 index 0000000000..639dc55e6c --- /dev/null +++ b/test/tint/bug/tint/1573.wgsl.expected.wgsl @@ -0,0 +1,7 @@ +@group(0) @binding(0) var a : atomic; + +@stage(compute) @workgroup_size(16) +fn main() { + var value = 42u; + let result = atomicCompareExchangeWeak(&(a), 0u, value); +}