tint: fix emitting duplicate structs for atomicCompareExchangeWeak
Bug: tint:1574
Change-Id: Id4ae2d2de9ac4678260f4ecfb3a0f779d170f9a4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/92280
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index eb44f14..e07139b 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -920,7 +920,7 @@
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<sem::Struct>())) {
+ if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -2822,6 +2822,14 @@
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 bcf84b0..819c79b 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -411,6 +411,12 @@
/// @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 @@
std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
std::unordered_map<const sem::Vector*, std::string> int_dot_funcs_;
std::unordered_map<const sem::Type*, std::string> float_modulo_funcs_;
+ std::unordered_set<const sem::Struct*> 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 137c466..19af4fa 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -1767,7 +1767,7 @@
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<sem::Struct>())) {
+ if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
return false;
}
@@ -3921,6 +3921,14 @@
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 c58d004..af7e4c9 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -411,6 +411,12 @@
/// @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 @@
std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_vector_write_;
std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_scalar_write_;
std::unordered_map<const sem::Type*, std::string> value_or_one_if_zero_;
+ std::unordered_set<const sem::Struct*> 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 cd984b2..3228e86 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -826,46 +826,66 @@
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<sem::Struct>())) {
- return false;
- }
-
auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As<sem::Pointer>();
auto sc = ptr_ty->StorageClass();
+ auto* str = builtin->ReturnType()->As<sem::Struct>();
- auto func = utils::GetOrCreate(atomicCompareExchangeWeak_, sc, [&]() -> std::string {
- auto name = UniqueIdentifier("atomicCompareExchangeWeak");
- auto& buf = helpers_;
-
- line(&buf) << "template <typename A, typename T>";
- {
- auto f = line(&buf);
- auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
- 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<sem::Struct>())) {
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<sem::Struct>());
+ 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 @@
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 21dee28..be98a86 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 <string>
+#include <tuple>
#include <unordered_map>
#include <unordered_set>
#include <vector>
@@ -332,6 +333,12 @@
/// @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 @@
/// type.
SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
- using StorageClassToString = std::unordered_map<ast::StorageClass, std::string>;
-
std::function<bool()> emit_continuing_;
/// Name of atomicCompareExchangeWeak() helper for the given pointer storage
- /// class.
- StorageClassToString atomicCompareExchangeWeak_;
+ /// class and struct return type
+ using ACEWKeyType =
+ utils::UnorderedKeyWrapper<std::tuple<ast::StorageClass, const sem::Struct*>>;
+ std::unordered_map<ACEWKeyType, std::string> atomicCompareExchangeWeak_;
/// Unique name of the 'TINT_INVARIANT' preprocessor define. Non-empty only if
/// an invariant attribute has been generated.
@@ -423,6 +430,7 @@
std::unordered_map<const sem::Builtin*, std::string> builtins_;
std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
std::unordered_map<uint32_t, std::string> int_dot_funcs_;
+ std::unordered_set<const sem::Struct*> 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 c284e1a..0aef172 100644
--- a/test/tint/bug/tint/1573.wgsl.expected.msl
+++ b/test/tint/bug/tint/1573.wgsl.expected.msl
@@ -6,9 +6,8 @@
uint old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 0000000..3a52312
--- /dev/null
+++ b/test/tint/bug/tint/1574.wgsl
@@ -0,0 +1,39 @@
+@group(0) @binding(0)
+var<storage, read_write> a_u32 : atomic<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> a_i32 : atomic<i32>;
+
+var<workgroup> b_u32 : atomic<u32>;
+
+var<workgroup> b_i32 : atomic<i32>;
+
+
+@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 0000000..bf711e7
--- /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 0000000..aad4961
--- /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 0000000..80ab0b6
--- /dev/null
+++ b/test/tint/bug/tint/1574.wgsl.expected.msl
@@ -0,0 +1,75 @@
+#include <metal_stdlib>
+
+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 0000000..be3e7f0
--- /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 0000000..97dbbee
--- /dev/null
+++ b/test/tint/bug/tint/1574.wgsl.expected.wgsl
@@ -0,0 +1,35 @@
+@group(0) @binding(0) var<storage, read_write> a_u32 : atomic<u32>;
+
+@group(0) @binding(1) var<storage, read_write> a_i32 : atomic<i32>;
+
+var<workgroup> b_u32 : atomic<u32>;
+
+var<workgroup> b_i32 : atomic<i32>;
+
+@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 6cab275..9f0ca2d 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 @@
int old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 94166ca..9437613 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 @@
uint old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 57f5ce3..effc141 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 @@
uint old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 0d754bc..2843d71 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 @@
int old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 11d8177..d2b6691 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 @@
int old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 545803b..f4fda56 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 @@
uint old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 6e181e3..811dd26 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 @@
uint old_value;
bool exchanged;
};
-template <typename A, typename T>
-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 38f9dc5..de5d73f 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 @@
int old_value;
bool exchanged;
};
-template <typename A, typename T>
-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};
}
