diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc index eb44f14c32..e07139bc8e 100644 --- a/src/tint/writer/glsl/generator_impl.cc +++ b/src/tint/writer/glsl/generator_impl.cc @@ -920,7 +920,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, case sem::BuiltinType::kAtomicCompareExchangeWeak: { // Emit the builtin return type unique to this overload. This does not // exist in the AST, so it will not be generated in Generate(). - if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { + if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As())) { return false; } @@ -2822,6 +2822,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { return true; } +bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) { + auto it = emitted_structs_.emplace(str); + if (!it.second) { + return true; + } + return EmitStructType(buffer, str); +} + bool GeneratorImpl::EmitStructMembers(TextBuffer* b, const sem::Struct* str, bool emit_offsets) { ScopedIndent si(b); for (auto* mem : str->Members()) { diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h index bcf84b09ef..819c79b973 100644 --- a/src/tint/writer/glsl/generator_impl.h +++ b/src/tint/writer/glsl/generator_impl.h @@ -411,6 +411,12 @@ class GeneratorImpl : public TextGenerator { /// @param ty the struct to generate /// @returns true if the struct is emitted bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty); + /// Handles generating a structure declaration only the first time called. Subsequent calls are + /// a no-op and return true. + /// @param buffer the text buffer that the type declaration will be written to + /// @param ty the struct to generate + /// @returns true if the struct is emitted + bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty); /// Handles generating the members of a structure /// @param buffer the text buffer that the struct members will be written to /// @param ty the struct to generate @@ -503,6 +509,7 @@ class GeneratorImpl : public TextGenerator { std::unordered_map dynamic_vector_write_; std::unordered_map int_dot_funcs_; std::unordered_map float_modulo_funcs_; + std::unordered_set emitted_structs_; bool requires_oes_sample_variables_ = false; bool requires_default_precision_qualifier_ = false; Version version_; diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 137c4667dd..19af4fad52 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -1767,7 +1767,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, case sem::BuiltinType::kAtomicCompareExchangeWeak: { // Emit the builtin return type unique to this overload. This does not // exist in the AST, so it will not be generated in Generate(). - if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { + if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As())) { return false; } @@ -3921,6 +3921,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { return true; } +bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) { + auto it = emitted_structs_.emplace(str); + if (!it.second) { + return true; + } + return EmitStructType(buffer, str); +} + bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) { switch (expr->op) { case ast::UnaryOp::kIndirection: diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h index c58d004edf..af7e4c9820 100644 --- a/src/tint/writer/hlsl/generator_impl.h +++ b/src/tint/writer/hlsl/generator_impl.h @@ -411,6 +411,12 @@ class GeneratorImpl : public TextGenerator { /// @param ty the struct to generate /// @returns true if the struct is emitted bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty); + /// Handles generating a structure declaration only the first time called. Subsequent calls are + /// a no-op and return true. + /// @param buffer the text buffer that the type declaration will be written to + /// @param ty the struct to generate + /// @returns true if the struct is emitted + bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty); /// Handles a unary op expression /// @param out the output of the expression stream /// @param expr the expression to emit @@ -530,6 +536,7 @@ class GeneratorImpl : public TextGenerator { std::unordered_map dynamic_matrix_vector_write_; std::unordered_map dynamic_matrix_scalar_write_; std::unordered_map value_or_one_if_zero_; + std::unordered_set emitted_structs_; }; } // namespace tint::writer::hlsl diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index cd984b26bb..3228e867d3 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -826,46 +826,66 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out, return call("atomic_exchange_explicit", true); case sem::BuiltinType::kAtomicCompareExchangeWeak: { - // Emit the builtin return type unique to this overload. This does not - // exist in the AST, so it will not be generated in Generate(). - if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { - return false; - } - auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As(); auto sc = ptr_ty->StorageClass(); + auto* str = builtin->ReturnType()->As(); - auto func = utils::GetOrCreate(atomicCompareExchangeWeak_, sc, [&]() -> std::string { - auto name = UniqueIdentifier("atomicCompareExchangeWeak"); - auto& buf = helpers_; - - line(&buf) << "template "; - { - auto f = line(&buf); - auto str_name = StructName(builtin->ReturnType()->As()); - f << str_name << " " << name << "("; - if (!EmitStorageClass(f, sc)) { + auto func = utils::GetOrCreate( + atomicCompareExchangeWeak_, ACEWKeyType{{sc, str}}, [&]() -> std::string { + // Emit the builtin return type unique to this overload. This does not + // exist in the AST, so it will not be generated in Generate(). + if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As())) { return ""; } - f << " A* atomic, T compare, T value) {"; - } - buf.IncrementIndent(); - TINT_DEFER({ - buf.DecrementIndent(); - line(&buf) << "}"; - line(&buf); + auto name = UniqueIdentifier("atomicCompareExchangeWeak"); + auto& buf = helpers_; + auto* atomic_ty = builtin->Parameters()[0]->Type(); + auto* arg_ty = builtin->Parameters()[1]->Type(); + + { + auto f = line(&buf); + auto str_name = StructName(builtin->ReturnType()->As()); + f << str_name << " " << name << "("; + if (!EmitTypeAndName(f, atomic_ty, "atomic")) { + return ""; + } + f << ", "; + if (!EmitTypeAndName(f, arg_ty, "compare")) { + return ""; + } + f << ", "; + if (!EmitTypeAndName(f, arg_ty, "value")) { + return ""; + } + f << ") {"; + } + + buf.IncrementIndent(); + TINT_DEFER({ + buf.DecrementIndent(); + line(&buf) << "}"; + line(&buf); + }); + + { + auto f = line(&buf); + if (!EmitTypeAndName(f, arg_ty, "old_value")) { + return ""; + } + f << " = compare;"; + } + line(&buf) << "bool exchanged = " + "atomic_compare_exchange_weak_explicit(atomic, " + "&old_value, value, memory_order_relaxed, " + "memory_order_relaxed);"; + line(&buf) << "return {old_value, exchanged};"; + return name; }); - line(&buf) << "T old_value = compare;"; - line(&buf) << "bool exchanged = " - "atomic_compare_exchange_weak_explicit(atomic, " - "&old_value, value, memory_order_relaxed, " - "memory_order_relaxed);"; - line(&buf) << "return {old_value, exchanged};"; - return name; - }); - + if (func.empty()) { + return false; + } return call(func, false); } @@ -2765,6 +2785,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { return true; } +bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) { + auto it = emitted_structs_.emplace(str); + if (!it.second) { + return true; + } + return EmitStructType(buffer, str); +} + bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) { // Handle `-e` when `e` is signed, so that we ensure that if `e` is the // largest negative value, it returns `e`. diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h index 21dee2848a..be98a86b43 100644 --- a/src/tint/writer/msl/generator_impl.h +++ b/src/tint/writer/msl/generator_impl.h @@ -16,6 +16,7 @@ #define SRC_TINT_WRITER_MSL_GENERATOR_IMPL_H_ #include +#include #include #include #include @@ -332,6 +333,12 @@ class GeneratorImpl : public TextGenerator { /// @param str the struct to generate /// @returns true if the struct is emitted bool EmitStructType(TextBuffer* buffer, const sem::Struct* str); + /// Handles generating a structure declaration only the first time called. Subsequent calls are + /// a no-op and return true. + /// @param buffer the text buffer that the type declaration will be written to + /// @param ty the struct to generate + /// @returns true if the struct is emitted + bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty); /// Handles a unary op expression /// @param out the output of the expression stream /// @param expr the expression to emit @@ -400,13 +407,13 @@ class GeneratorImpl : public TextGenerator { /// type. SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty); - using StorageClassToString = std::unordered_map; - std::function emit_continuing_; /// Name of atomicCompareExchangeWeak() helper for the given pointer storage - /// class. - StorageClassToString atomicCompareExchangeWeak_; + /// class and struct return type + using ACEWKeyType = + utils::UnorderedKeyWrapper>; + std::unordered_map atomicCompareExchangeWeak_; /// Unique name of the 'TINT_INVARIANT' preprocessor define. Non-empty only if /// an invariant attribute has been generated. @@ -423,6 +430,7 @@ class GeneratorImpl : public TextGenerator { std::unordered_map builtins_; std::unordered_map unary_minus_funcs_; std::unordered_map int_dot_funcs_; + std::unordered_set emitted_structs_; }; } // namespace tint::writer::msl diff --git a/test/tint/bug/tint/1573.wgsl.expected.msl b/test/tint/bug/tint/1573.wgsl.expected.msl index c284e1a5ea..0aef1728fa 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.msl +++ b/test/tint/bug/tint/1573.wgsl.expected.msl @@ -6,9 +6,8 @@ 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; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/bug/tint/1574.wgsl b/test/tint/bug/tint/1574.wgsl new file mode 100644 index 0000000000..3a52312c56 --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl @@ -0,0 +1,39 @@ +@group(0) @binding(0) +var a_u32 : atomic; + +@group(0) @binding(1) +var a_i32 : atomic; + +var b_u32 : atomic; + +var b_i32 : atomic; + + +@stage(compute) @workgroup_size(16) +fn main() { + { + var value = 42u; + let r1 = atomicCompareExchangeWeak(&a_u32, 0u, value); + let r2 = atomicCompareExchangeWeak(&a_u32, 0u, value); + let r3 = atomicCompareExchangeWeak(&a_u32, 0u, value); + } + { + var value = 42; + let r1 = atomicCompareExchangeWeak(&a_i32, 0, value); + let r2 = atomicCompareExchangeWeak(&a_i32, 0, value); + let r3 = atomicCompareExchangeWeak(&a_i32, 0, value); + } + { + var value = 42u; + let r1 = atomicCompareExchangeWeak(&b_u32, 0u, value); + let r2 = atomicCompareExchangeWeak(&b_u32, 0u, value); + let r3 = atomicCompareExchangeWeak(&b_u32, 0u, value); + } + { + var value = 42; + let r1 = atomicCompareExchangeWeak(&b_i32, 0, value); + let r2 = atomicCompareExchangeWeak(&b_i32, 0, value); + let r3 = atomicCompareExchangeWeak(&b_i32, 0, value); + } + +} diff --git a/test/tint/bug/tint/1574.wgsl.expected.glsl b/test/tint/bug/tint/1574.wgsl.expected.glsl new file mode 100644 index 0000000000..bf711e7891 --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl.expected.glsl @@ -0,0 +1,102 @@ +#version 310 es + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; + + +struct a_u32_block { + uint inner; +}; + +layout(binding = 0, std430) buffer a_u32_block_1 { + uint inner; +} a_u32; +struct a_i32_block { + int inner; +}; + +layout(binding = 1, std430) buffer a_i32_block_1 { + int inner; +} a_i32; +shared uint b_u32; +shared int b_i32; +void tint_symbol(uint local_invocation_index) { + if ((local_invocation_index < 1u)) { + atomicExchange(b_u32, 0u); + atomicExchange(b_i32, 0); + } + barrier(); + { + uint value = 42u; + atomic_compare_exchange_resultu32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(a_u32.inner, 0u, value); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u; + atomic_compare_exchange_resultu32 r1 = atomic_compare_result; + atomic_compare_exchange_resultu32 atomic_compare_result_1; + atomic_compare_result_1.old_value = atomicCompSwap(a_u32.inner, 0u, value); + atomic_compare_result_1.exchanged = atomic_compare_result_1.old_value == 0u; + atomic_compare_exchange_resultu32 r2 = atomic_compare_result_1; + atomic_compare_exchange_resultu32 atomic_compare_result_2; + atomic_compare_result_2.old_value = atomicCompSwap(a_u32.inner, 0u, value); + atomic_compare_result_2.exchanged = atomic_compare_result_2.old_value == 0u; + atomic_compare_exchange_resultu32 r3 = atomic_compare_result_2; + } + { + int value = 42; + atomic_compare_exchange_resulti32 atomic_compare_result_3; + atomic_compare_result_3.old_value = atomicCompSwap(a_i32.inner, 0, value); + atomic_compare_result_3.exchanged = atomic_compare_result_3.old_value == 0; + atomic_compare_exchange_resulti32 r1 = atomic_compare_result_3; + atomic_compare_exchange_resulti32 atomic_compare_result_4; + atomic_compare_result_4.old_value = atomicCompSwap(a_i32.inner, 0, value); + atomic_compare_result_4.exchanged = atomic_compare_result_4.old_value == 0; + atomic_compare_exchange_resulti32 r2 = atomic_compare_result_4; + atomic_compare_exchange_resulti32 atomic_compare_result_5; + atomic_compare_result_5.old_value = atomicCompSwap(a_i32.inner, 0, value); + atomic_compare_result_5.exchanged = atomic_compare_result_5.old_value == 0; + atomic_compare_exchange_resulti32 r3 = atomic_compare_result_5; + } + { + uint value = 42u; + atomic_compare_exchange_resultu32 atomic_compare_result_6; + atomic_compare_result_6.old_value = atomicCompSwap(b_u32, 0u, value); + atomic_compare_result_6.exchanged = atomic_compare_result_6.old_value == 0u; + atomic_compare_exchange_resultu32 r1 = atomic_compare_result_6; + atomic_compare_exchange_resultu32 atomic_compare_result_7; + atomic_compare_result_7.old_value = atomicCompSwap(b_u32, 0u, value); + atomic_compare_result_7.exchanged = atomic_compare_result_7.old_value == 0u; + atomic_compare_exchange_resultu32 r2 = atomic_compare_result_7; + atomic_compare_exchange_resultu32 atomic_compare_result_8; + atomic_compare_result_8.old_value = atomicCompSwap(b_u32, 0u, value); + atomic_compare_result_8.exchanged = atomic_compare_result_8.old_value == 0u; + atomic_compare_exchange_resultu32 r3 = atomic_compare_result_8; + } + { + int value = 42; + atomic_compare_exchange_resulti32 atomic_compare_result_9; + atomic_compare_result_9.old_value = atomicCompSwap(b_i32, 0, value); + atomic_compare_result_9.exchanged = atomic_compare_result_9.old_value == 0; + atomic_compare_exchange_resulti32 r1 = atomic_compare_result_9; + atomic_compare_exchange_resulti32 atomic_compare_result_10; + atomic_compare_result_10.old_value = atomicCompSwap(b_i32, 0, value); + atomic_compare_result_10.exchanged = atomic_compare_result_10.old_value == 0; + atomic_compare_exchange_resulti32 r2 = atomic_compare_result_10; + atomic_compare_exchange_resulti32 atomic_compare_result_11; + atomic_compare_result_11.old_value = atomicCompSwap(b_i32, 0, value); + atomic_compare_result_11.exchanged = atomic_compare_result_11.old_value == 0; + atomic_compare_exchange_resulti32 r3 = atomic_compare_result_11; + } +} + +layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(gl_LocalInvocationIndex); + return; +} diff --git a/test/tint/bug/tint/1574.wgsl.expected.hlsl b/test/tint/bug/tint/1574.wgsl.expected.hlsl new file mode 100644 index 0000000000..aad4961f4a --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl.expected.hlsl @@ -0,0 +1,105 @@ +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; +RWByteAddressBuffer a_u32 : register(u0, space0); +RWByteAddressBuffer a_i32 : register(u1, space0); +groupshared uint b_u32; +groupshared int b_i32; + +struct tint_symbol_1 { + uint local_invocation_index : SV_GroupIndex; +}; +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; +} + + +struct atomic_compare_exchange_weak_ret_type_1 { + int old_value; + bool exchanged; +}; + +atomic_compare_exchange_weak_ret_type_1 tint_atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, int compare, int value) { + atomic_compare_exchange_weak_ret_type_1 result=(atomic_compare_exchange_weak_ret_type_1)0; + buffer.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +void main_inner(uint local_invocation_index) { + if ((local_invocation_index < 1u)) { + uint atomic_result = 0u; + InterlockedExchange(b_u32, 0u, atomic_result); + int atomic_result_1 = 0; + InterlockedExchange(b_i32, 0, atomic_result_1); + } + GroupMemoryBarrierWithGroupSync(); + { + uint value = 42u; + const atomic_compare_exchange_weak_ret_type r1 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value); + const atomic_compare_exchange_weak_ret_type r2 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value); + const atomic_compare_exchange_weak_ret_type r3 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value); + } + { + int value = 42; + const atomic_compare_exchange_weak_ret_type_1 r1 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value); + const atomic_compare_exchange_weak_ret_type_1 r2 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value); + const atomic_compare_exchange_weak_ret_type_1 r3 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value); + } + { + uint value = 42u; + atomic_compare_exchange_resultu32 atomic_result_2 = (atomic_compare_exchange_resultu32)0; + uint atomic_compare_value = 0u; + InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value); + atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value; + const atomic_compare_exchange_resultu32 r1 = atomic_result_2; + atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0; + uint atomic_compare_value_1 = 0u; + InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value); + atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1; + const atomic_compare_exchange_resultu32 r2 = atomic_result_3; + atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0; + uint atomic_compare_value_2 = 0u; + InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value); + atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2; + const atomic_compare_exchange_resultu32 r3 = atomic_result_4; + } + { + int value = 42; + atomic_compare_exchange_resulti32 atomic_result_5 = (atomic_compare_exchange_resulti32)0; + int atomic_compare_value_3 = 0; + InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value); + atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3; + const atomic_compare_exchange_resulti32 r1 = atomic_result_5; + atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0; + int atomic_compare_value_4 = 0; + InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value); + atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4; + const atomic_compare_exchange_resulti32 r2 = atomic_result_6; + atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0; + int atomic_compare_value_5 = 0; + InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value); + atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5; + const atomic_compare_exchange_resulti32 r3 = atomic_result_7; + } +} + +[numthreads(16, 1, 1)] +void main(tint_symbol_1 tint_symbol) { + main_inner(tint_symbol.local_invocation_index); + return; +} diff --git a/test/tint/bug/tint/1574.wgsl.expected.msl b/test/tint/bug/tint/1574.wgsl.expected.msl new file mode 100644 index 0000000000..80ab0b68be --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl.expected.msl @@ -0,0 +1,75 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_2(device atomic_int* atomic, int compare, int value) { + int old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_3(threadgroup atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_4(threadgroup atomic_int* atomic, int compare, int value) { + int old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +void tint_symbol_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1, threadgroup atomic_int* const tint_symbol_2, device atomic_uint* const tint_symbol_3, device atomic_int* const tint_symbol_4) { + if ((local_invocation_index < 1u)) { + atomic_store_explicit(tint_symbol_1, 0u, memory_order_relaxed); + atomic_store_explicit(tint_symbol_2, 0, memory_order_relaxed); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + { + uint value = 42u; + atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + } + { + int value = 42; + atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + } + { + uint value = 42u; + atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + } + { + int value = 42; + atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + } +} + +kernel void tint_symbol(device atomic_uint* tint_symbol_7 [[buffer(0)]], device atomic_int* tint_symbol_8 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) { + threadgroup atomic_uint tint_symbol_5; + threadgroup atomic_int tint_symbol_6; + tint_symbol_inner(local_invocation_index, &(tint_symbol_5), &(tint_symbol_6), tint_symbol_7, tint_symbol_8); + return; +} + diff --git a/test/tint/bug/tint/1574.wgsl.expected.spvasm b/test/tint/bug/tint/1574.wgsl.expected.spvasm new file mode 100644 index 0000000000..be3e7f0a6a --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl.expected.spvasm @@ -0,0 +1,158 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 118 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %local_invocation_index_1 + OpExecutionMode %main LocalSize 16 1 1 + OpName %local_invocation_index_1 "local_invocation_index_1" + OpName %a_u32_block "a_u32_block" + OpMemberName %a_u32_block 0 "inner" + OpName %a_u32 "a_u32" + OpName %a_i32_block "a_i32_block" + OpMemberName %a_i32_block 0 "inner" + OpName %a_i32 "a_i32" + OpName %b_u32 "b_u32" + OpName %b_i32 "b_i32" + OpName %main_inner "main_inner" + OpName %local_invocation_index "local_invocation_index" + 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" + OpName %value_0 "value" + OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" + OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %value_1 "value" + OpName %value_2 "value" + OpName %main "main" + OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex + OpDecorate %a_u32_block Block + OpMemberDecorate %a_u32_block 0 Offset 0 + OpDecorate %a_u32 DescriptorSet 0 + OpDecorate %a_u32 Binding 0 + OpDecorate %a_i32_block Block + OpMemberDecorate %a_i32_block 0 Offset 0 + OpDecorate %a_i32 DescriptorSet 0 + OpDecorate %a_i32 Binding 1 + OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + %uint = OpTypeInt 32 0 +%_ptr_Input_uint = OpTypePointer Input %uint +%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input +%a_u32_block = OpTypeStruct %uint +%_ptr_StorageBuffer_a_u32_block = OpTypePointer StorageBuffer %a_u32_block + %a_u32 = OpVariable %_ptr_StorageBuffer_a_u32_block StorageBuffer + %int = OpTypeInt 32 1 +%a_i32_block = OpTypeStruct %int +%_ptr_StorageBuffer_a_i32_block = OpTypePointer StorageBuffer %a_i32_block + %a_i32 = OpVariable %_ptr_StorageBuffer_a_i32_block StorageBuffer +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %b_u32 = OpVariable %_ptr_Workgroup_uint Workgroup +%_ptr_Workgroup_int = OpTypePointer Workgroup %int + %b_i32 = OpVariable %_ptr_Workgroup_int Workgroup + %void = OpTypeVoid + %15 = OpTypeFunction %void %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %29 = OpConstantNull %uint + %32 = OpConstantNull %int + %uint_264 = OpConstant %uint 264 + %uint_42 = OpConstant %uint 42 +%_ptr_Function_uint = OpTypePointer Function %uint +%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %int_42 = OpConstant %int 42 +%_ptr_Function_int = OpTypePointer Function %int +%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int + %113 = OpTypeFunction %void + %main_inner = OpFunction %void None %15 +%local_invocation_index = OpFunctionParameter %uint + %19 = OpLabel + %value = OpVariable %_ptr_Function_uint Function %29 + %value_0 = OpVariable %_ptr_Function_int Function %32 + %value_1 = OpVariable %_ptr_Function_uint Function %29 + %value_2 = OpVariable %_ptr_Function_int Function %32 + %21 = OpULessThan %bool %local_invocation_index %uint_1 + OpSelectionMerge %23 None + OpBranchConditional %21 %24 %23 + %24 = OpLabel + OpAtomicStore %b_u32 %uint_2 %uint_0 %29 + OpAtomicStore %b_i32 %uint_2 %uint_0 %32 + OpBranch %23 + %23 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + OpStore %value %uint_42 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0 + %43 = OpLoad %uint %value + %44 = OpAtomicCompareExchange %uint %42 %uint_1 %uint_0 %uint_0 %43 %29 + %45 = OpIEqual %bool %44 %43 + %38 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %44 %45 + %48 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0 + %49 = OpLoad %uint %value + %50 = OpAtomicCompareExchange %uint %48 %uint_1 %uint_0 %uint_0 %49 %29 + %51 = OpIEqual %bool %50 %49 + %46 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %50 %51 + %54 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0 + %55 = OpLoad %uint %value + %56 = OpAtomicCompareExchange %uint %54 %uint_1 %uint_0 %uint_0 %55 %29 + %57 = OpIEqual %bool %56 %55 + %52 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %56 %57 + OpStore %value_0 %int_42 + %65 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 + %66 = OpLoad %int %value_0 + %67 = OpAtomicCompareExchange %int %65 %uint_1 %uint_0 %uint_0 %66 %32 + %68 = OpIEqual %bool %67 %66 + %61 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %67 %68 + %71 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 + %72 = OpLoad %int %value_0 + %73 = OpAtomicCompareExchange %int %71 %uint_1 %uint_0 %uint_0 %72 %32 + %74 = OpIEqual %bool %73 %72 + %69 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %73 %74 + %77 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 + %78 = OpLoad %int %value_0 + %79 = OpAtomicCompareExchange %int %77 %uint_1 %uint_0 %uint_0 %78 %32 + %80 = OpIEqual %bool %79 %78 + %75 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %79 %80 + OpStore %value_1 %uint_42 + %84 = OpLoad %uint %value_1 + %85 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %84 %29 + %86 = OpIEqual %bool %85 %84 + %82 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %85 %86 + %89 = OpLoad %uint %value_1 + %90 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %89 %29 + %91 = OpIEqual %bool %90 %89 + %87 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %90 %91 + %94 = OpLoad %uint %value_1 + %95 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %94 %29 + %96 = OpIEqual %bool %95 %94 + %92 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %95 %96 + OpStore %value_2 %int_42 + %100 = OpLoad %int %value_2 + %101 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %100 %32 + %102 = OpIEqual %bool %101 %100 + %98 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %101 %102 + %105 = OpLoad %int %value_2 + %106 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %105 %32 + %107 = OpIEqual %bool %106 %105 + %103 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %106 %107 + %110 = OpLoad %int %value_2 + %111 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %110 %32 + %112 = OpIEqual %bool %111 %110 + %108 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %111 %112 + OpReturn + OpFunctionEnd + %main = OpFunction %void None %113 + %115 = OpLabel + %117 = OpLoad %uint %local_invocation_index_1 + %116 = OpFunctionCall %void %main_inner %117 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1574.wgsl.expected.wgsl b/test/tint/bug/tint/1574.wgsl.expected.wgsl new file mode 100644 index 0000000000..97dbbee365 --- /dev/null +++ b/test/tint/bug/tint/1574.wgsl.expected.wgsl @@ -0,0 +1,35 @@ +@group(0) @binding(0) var a_u32 : atomic; + +@group(0) @binding(1) var a_i32 : atomic; + +var b_u32 : atomic; + +var b_i32 : atomic; + +@stage(compute) @workgroup_size(16) +fn main() { + { + var value = 42u; + let r1 = atomicCompareExchangeWeak(&(a_u32), 0u, value); + let r2 = atomicCompareExchangeWeak(&(a_u32), 0u, value); + let r3 = atomicCompareExchangeWeak(&(a_u32), 0u, value); + } + { + var value = 42; + let r1 = atomicCompareExchangeWeak(&(a_i32), 0, value); + let r2 = atomicCompareExchangeWeak(&(a_i32), 0, value); + let r3 = atomicCompareExchangeWeak(&(a_i32), 0, value); + } + { + var value = 42u; + let r1 = atomicCompareExchangeWeak(&(b_u32), 0u, value); + let r2 = atomicCompareExchangeWeak(&(b_u32), 0u, value); + let r3 = atomicCompareExchangeWeak(&(b_u32), 0u, value); + } + { + var value = 42; + let r1 = atomicCompareExchangeWeak(&(b_i32), 0, value); + let r2 = atomicCompareExchangeWeak(&(b_i32), 0, value); + let r3 = atomicCompareExchangeWeak(&(b_i32), 0, value); + } +} diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl index 6cab275aa3..9f0ca2d9ad 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 { int old_value; bool exchanged; }; -template -atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { + int old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl index 94166ca79b..9437613622 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl @@ -6,9 +6,8 @@ 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; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl index 57f5ce34d3..effc1417eb 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; }; -template -atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl index 0d754bca1b..2843d71cf2 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 { int old_value; bool exchanged; }; -template -atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { + int old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl index 11d8177555..d2b66917f6 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 { int old_value; bool exchanged; }; -template -atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { + int old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl index 545803b2ee..f4fda565e0 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl @@ -6,9 +6,8 @@ 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; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl index 6e181e37d5..811dd26244 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 { uint old_value; bool exchanged; }; -template -atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { + uint old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl index 38f9dc57d2..de5d73f584 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl @@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 { int old_value; bool exchanged; }; -template -atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T old_value = compare; +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { + int old_value = compare; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; }