writer/msl: Implement atomics
Common logic between the HLSL, WGSL and MSL writers has been moved into
the TextGenerator base class.
Fixed: tint:892
Change-Id: I0f469516947fe64817ce6251e436da74e5e176e8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56068
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 14a78f9..c808b43 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -669,6 +669,7 @@
utils/unique_vector_test.cc
writer/append_vector_test.cc
writer/float_to_string_test.cc
+ writer/text_generator_test.cc
)
if(${TINT_BUILD_SPV_READER})
diff --git a/src/resolver/atomics_validation_test.cc b/src/resolver/atomics_validation_test.cc
index 767a648..4788d45 100644
--- a/src/resolver/atomics_validation_test.cc
+++ b/src/resolver/atomics_validation_test.cc
@@ -61,7 +61,7 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: cannot declare an atomic var in a function scope");
+ "12:34 error: atomic var requires workgroup storage");
}
TEST_F(ResolverAtomicValidationTest, NoAtomicExpr) {
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index c6462df..f613790 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -909,21 +909,13 @@
// https://gpuweb.github.io/gpuweb/wgsl/#atomic-types
// Atomic types may only be instantiated by variables in the workgroup storage
// class or by storage buffer variables with a read_write access mode.
- if (info->type->UnwrapRef()->Is<sem::Atomic>()) {
- if (info->kind != VariableKind::kGlobal) {
- // Neither storage nor workgroup storage classes can be used in function
- // scopes.
- AddError("cannot declare an atomic var in a function scope",
- info->declaration->type()->source());
- return false;
- }
- if (info->storage_class != ast::StorageClass::kWorkgroup) {
- // Storage buffers require a structure, so just check for workgroup
- // storage here.
- AddError("atomic var requires workgroup storage",
- info->declaration->type()->source());
- return false;
- }
+ if (info->type->UnwrapRef()->Is<sem::Atomic>() &&
+ info->storage_class != ast::StorageClass::kWorkgroup) {
+ // Storage buffers require a structure, so just check for workgroup
+ // storage here.
+ AddError("atomic var requires workgroup storage",
+ info->declaration->type()->source());
+ return false;
}
return true;
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 430f05f..9e81adb 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -104,8 +104,7 @@
} // namespace
-GeneratorImpl::GeneratorImpl(const Program* program)
- : builder_(ProgramBuilder::Wrap(program)) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
GeneratorImpl::~GeneratorImpl() = default;
@@ -165,10 +164,6 @@
return true;
}
-std::string GeneratorImpl::generate_name(const std::string& prefix) {
- return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
-}
-
bool GeneratorImpl::EmitArrayAccessor(std::ostream& out,
ast::ArrayAccessorExpression* expr) {
if (!EmitExpression(out, expr->array())) {
@@ -222,7 +217,7 @@
bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
if (expr->op() == ast::BinaryOp::kLogicalAnd ||
expr->op() == ast::BinaryOp::kLogicalOr) {
- auto name = generate_name(kTempNamePrefix);
+ auto name = UniqueIdentifier(kTempNamePrefix);
{
auto pre = line();
@@ -505,7 +500,7 @@
const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
const auto& params = expr->params();
- std::string scalar_offset = generate_name("scalar_offset");
+ std::string scalar_offset = UniqueIdentifier("scalar_offset");
{
auto pre = line();
pre << "const int " << scalar_offset << " = (";
@@ -534,7 +529,7 @@
};
// Has a minimum alignment of 8 bytes, so is either .xy or .zw
auto load_vec2 = [&] {
- std::string ubo_load = generate_name("ubo_load");
+ std::string ubo_load = UniqueIdentifier("ubo_load");
{
auto pre = line();
@@ -744,7 +739,7 @@
transform::DecomposeMemoryAccess::Intrinsic::Op op) {
using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
- std::string result = generate_name("atomic_result");
+ std::string result = UniqueIdentifier("atomic_result");
auto* result_ty = TypeOf(expr);
if (!result_ty->Is<sem::Void>()) {
@@ -849,7 +844,7 @@
auto* compare_value = expr->params()[2];
auto* value = expr->params()[3];
- std::string compare = generate_name("atomic_compare_value");
+ std::string compare = UniqueIdentifier("atomic_compare_value");
{ // T atomic_compare_value = compare_value;
auto pre = line();
if (!EmitTypeAndName(pre, TypeOf(compare_value),
@@ -924,7 +919,7 @@
bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
- std::string result = generate_name("atomic_result");
+ std::string result = UniqueIdentifier("atomic_result");
if (!intrinsic->ReturnType()->Is<sem::Void>()) {
auto pre = line();
@@ -1018,7 +1013,7 @@
auto* compare_value = expr->params()[1];
auto* value = expr->params()[2];
- std::string compare = generate_name("atomic_compare_value");
+ std::string compare = UniqueIdentifier("atomic_compare_value");
{ // T compare_value = <compare_value>;
auto pre = line();
@@ -1130,8 +1125,8 @@
// Exponent is an integer, which HLSL does not have an overload for.
// We need to cast from a float.
- auto float_exp = generate_name(kTempNamePrefix);
- auto significand = generate_name(kTempNamePrefix);
+ auto float_exp = UniqueIdentifier(kTempNamePrefix);
+ auto significand = UniqueIdentifier(kTempNamePrefix);
line() << "float" << width << " " << float_exp << ";";
{
auto pre = line();
@@ -1173,8 +1168,8 @@
constexpr auto* kMinNormalExponent = "0x0080000";
constexpr auto* kMaxNormalExponent = "0x7f00000";
- auto exponent = generate_name("tint_isnormal_exponent");
- auto clamped = generate_name("tint_isnormal_clamped");
+ auto exponent = UniqueIdentifier("tint_isnormal_exponent");
+ auto clamped = UniqueIdentifier("tint_isnormal_clamped");
{
auto pre = line();
@@ -1196,7 +1191,7 @@
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
auto* param = expr->params()[0];
- auto tmp_name = generate_name(kTempNamePrefix);
+ auto tmp_name = UniqueIdentifier(kTempNamePrefix);
std::ostringstream expr_out;
if (!EmitExpression(expr_out, param)) {
return false;
@@ -1261,7 +1256,7 @@
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
auto* param = expr->params()[0];
- auto tmp_name = generate_name(kTempNamePrefix);
+ auto tmp_name = UniqueIdentifier(kTempNamePrefix);
std::ostringstream expr_out;
if (!EmitExpression(expr_out, param)) {
return false;
@@ -1282,7 +1277,7 @@
switch (intrinsic->Type()) {
case sem::IntrinsicType::kUnpack4x8snorm:
case sem::IntrinsicType::kUnpack2x16snorm: {
- auto tmp_name2 = generate_name(kTempNamePrefix);
+ auto tmp_name2 = UniqueIdentifier(kTempNamePrefix);
line() << "int " << tmp_name2 << " = int(" << expr_out.str() << ");";
{ // Perform sign extension on the converted values.
auto pre = line();
@@ -1302,7 +1297,7 @@
}
case sem::IntrinsicType::kUnpack4x8unorm:
case sem::IntrinsicType::kUnpack2x16unorm: {
- auto tmp_name2 = generate_name(kTempNamePrefix);
+ auto tmp_name2 = UniqueIdentifier(kTempNamePrefix);
line() << "uint " << tmp_name2 << " = " << expr_out.str() << ";";
{
auto pre = line();
@@ -1492,7 +1487,7 @@
}
// Declare a variable to hold the queried texture info
- auto dims = generate_name(kTempNamePrefix);
+ auto dims = UniqueIdentifier(kTempNamePrefix);
if (num_dimensions == 1) {
line() << "int " << dims << ";";
} else {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 77f2bcf..58b791d 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -361,11 +361,6 @@
ast::InterpolationType type,
ast::InterpolationSampling sampling) const;
- /// Generate a unique name
- /// @param prefix the name prefix
- /// @returns a unique name
- std::string generate_name(const std::string& prefix);
-
private:
enum class VarType { kIn, kOut };
@@ -376,25 +371,6 @@
std::string get_buffer_name(ast::Expression* expr);
- /// @returns the resolved type of the ast::Expression `expr`
- /// @param expr the expression
- sem::Type* TypeOf(ast::Expression* expr) const {
- return builder_.TypeOf(expr);
- }
-
- /// @returns the resolved type of the ast::Type `type`
- /// @param type the type
- const sem::Type* TypeOf(const ast::Type* type) const {
- return builder_.TypeOf(type);
- }
-
- /// @returns the resolved type of the ast::TypeDecl `type_decl`
- /// @param type_decl the type
- const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
- return builder_.TypeOf(type_decl);
- }
-
- ProgramBuilder builder_;
std::function<bool()> emit_continuing_;
std::unordered_map<const sem::Struct*, std::string> structure_builders_;
};
diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc
index 30cf659..cd01730 100644
--- a/src/writer/hlsl/generator_impl_test.cc
+++ b/src/writer/hlsl/generator_impl_test.cc
@@ -43,24 +43,6 @@
)");
}
-TEST_F(HlslGeneratorImplTest, InputStructName) {
- GeneratorImpl& gen = Build();
-
- ASSERT_EQ(gen.generate_name("func_main_in"), "func_main_in");
-}
-
-TEST_F(HlslGeneratorImplTest, InputStructName_ConflictWithExisting) {
- Symbols().Register("func_main_out_1");
- Symbols().Register("func_main_out_2");
-
- GeneratorImpl& gen = Build();
-
- ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out");
- ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_3");
- ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_4");
- ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_5");
-}
-
struct HlslBuiltinData {
ast::Builtin builtin;
const char* attribute_name;
diff --git a/src/writer/msl/generator.cc b/src/writer/msl/generator.cc
index 9f08f56..3c2190f 100644
--- a/src/writer/msl/generator.cc
+++ b/src/writer/msl/generator.cc
@@ -13,6 +13,7 @@
// limitations under the License.
#include "src/writer/msl/generator.h"
+#include "src/writer/msl/generator_impl.h"
namespace tint {
namespace writer {
diff --git a/src/writer/msl/generator.h b/src/writer/msl/generator.h
index 3844875..d63f755 100644
--- a/src/writer/msl/generator.h
+++ b/src/writer/msl/generator.h
@@ -18,13 +18,14 @@
#include <memory>
#include <string>
-#include "src/writer/msl/generator_impl.h"
#include "src/writer/text.h"
namespace tint {
namespace writer {
namespace msl {
+class GeneratorImpl;
+
/// Class to generate MSL source
class Generator : public Text {
public:
@@ -46,6 +47,9 @@
std::string error() const;
private:
+ Generator(const Generator&) = delete;
+ Generator& operator=(const Generator&) = delete;
+
std::unique_ptr<GeneratorImpl> impl_;
};
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index a12c44d..bf7ffab 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -33,6 +33,7 @@
#include "src/ast/variable_decl_statement.h"
#include "src/ast/void.h"
#include "src/sem/array.h"
+#include "src/sem/atomic_type.h"
#include "src/sem/bool_type.h"
#include "src/sem/call.h"
#include "src/sem/depth_texture_type.h"
@@ -71,8 +72,7 @@
} // namespace
-GeneratorImpl::GeneratorImpl(const Program* program)
- : TextGenerator(), program_(program) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
GeneratorImpl::~GeneratorImpl() = default;
@@ -359,6 +359,9 @@
bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
+ if (intrinsic->IsAtomic()) {
+ return EmitAtomicCall(out, expr, intrinsic);
+ }
if (intrinsic->IsTexture()) {
return EmitTextureCall(out, expr, intrinsic);
}
@@ -422,6 +425,111 @@
return true;
}
+bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
+ ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic) {
+ auto call = [&](const char* name) {
+ out << name;
+ {
+ ScopedParen sp(out);
+ for (size_t i = 0; i < expr->params().size(); i++) {
+ auto* arg = expr->params()[i];
+ if (i > 0) {
+ out << ", ";
+ }
+ if (!EmitExpression(out, arg)) {
+ return false;
+ }
+ }
+ out << ", memory_order_relaxed";
+ }
+ return true;
+ };
+
+ switch (intrinsic->Type()) {
+ case sem::IntrinsicType::kAtomicLoad:
+ return call("atomic_load_explicit");
+
+ case sem::IntrinsicType::kAtomicStore:
+ return call("atomic_store_explicit");
+
+ case sem::IntrinsicType::kAtomicAdd:
+ return call("atomic_fetch_add_explicit");
+
+ case sem::IntrinsicType::kAtomicMax:
+ return call("atomic_fetch_max_explicit");
+
+ case sem::IntrinsicType::kAtomicMin:
+ return call("atomic_fetch_min_explicit");
+
+ case sem::IntrinsicType::kAtomicAnd:
+ return call("atomic_fetch_and_explicit");
+
+ case sem::IntrinsicType::kAtomicOr:
+ return call("atomic_fetch_or_explicit");
+
+ case sem::IntrinsicType::kAtomicXor:
+ return call("atomic_fetch_xor_explicit");
+
+ case sem::IntrinsicType::kAtomicExchange:
+ return call("atomic_exchange_explicit");
+
+ case sem::IntrinsicType::kAtomicCompareExchangeWeak: {
+ auto* target = expr->params()[0];
+ auto* compare_value = expr->params()[1];
+ auto* value = expr->params()[2];
+
+ auto prev_value = UniqueIdentifier("prev_value");
+ auto matched = UniqueIdentifier("matched");
+
+ { // prev_value = <compare_value>;
+ auto pre = line();
+ if (!EmitType(pre, TypeOf(value), "")) {
+ return false;
+ }
+ pre << " " << prev_value << " = ";
+ if (!EmitExpression(pre, compare_value)) {
+ return false;
+ }
+ pre << ";";
+ }
+
+ { // bool matched = atomic_compare_exchange_weak_explicit(
+ // target, &got, <value>, memory_order_relaxed, memory_order_relaxed)
+ auto pre = line();
+ pre << "bool " << matched << " = atomic_compare_exchange_weak_explicit";
+ {
+ ScopedParen sp(pre);
+ if (!EmitExpression(pre, target)) {
+ return false;
+ }
+ pre << ", &" << prev_value << ", ";
+ if (!EmitExpression(pre, value)) {
+ return false;
+ }
+ pre << ", memory_order_relaxed, memory_order_relaxed";
+ }
+ pre << ";";
+ }
+
+ { // [u]int2(got, matched)
+ if (!EmitType(out, TypeOf(expr), "")) {
+ return false;
+ }
+ out << "(" << prev_value << ", " << matched << ")";
+ }
+ return true;
+ }
+
+ default:
+ break;
+ }
+
+ TINT_UNREACHABLE(Writer, diagnostics_)
+ << "unsupported atomic intrinsic: " << intrinsic->Type();
+ return false;
+}
+
bool GeneratorImpl::EmitTextureCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic) {
@@ -1550,6 +1658,20 @@
bool GeneratorImpl::EmitType(std::ostream& out,
const sem::Type* type,
const std::string& name) {
+ if (auto* atomic = type->As<sem::Atomic>()) {
+ if (atomic->Type()->Is<sem::I32>()) {
+ out << "atomic_int";
+ return true;
+ }
+ if (atomic->Type()->Is<sem::U32>()) {
+ out << "atomic_uint";
+ return true;
+ }
+ TINT_ICE(Writer, diagnostics_)
+ << "unhandled atomic type " << atomic->Type()->type_name();
+ return false;
+ }
+
if (auto* ary = type->As<sem::Array>()) {
const sem::Type* base_type = ary;
std::vector<uint32_t> sizes;
@@ -1570,18 +1692,33 @@
for (uint32_t size : sizes) {
out << "[" << size << "]";
}
- } else if (type->Is<sem::Bool>()) {
+ return true;
+ }
+
+ if (type->Is<sem::Bool>()) {
out << "bool";
- } else if (type->Is<sem::F32>()) {
+ return true;
+ }
+
+ if (type->Is<sem::F32>()) {
out << "float";
- } else if (type->Is<sem::I32>()) {
+ return true;
+ }
+
+ if (type->Is<sem::I32>()) {
out << "int";
- } else if (auto* mat = type->As<sem::Matrix>()) {
+ return true;
+ }
+
+ if (auto* mat = type->As<sem::Matrix>()) {
if (!EmitType(out, mat->type(), "")) {
return false;
}
out << mat->columns() << "x" << mat->rows();
- } else if (auto* ptr = type->As<sem::Pointer>()) {
+ return true;
+ }
+
+ if (auto* ptr = type->As<sem::Pointer>()) {
switch (ptr->StorageClass()) {
case ast::StorageClass::kFunction:
case ast::StorageClass::kPrivate:
@@ -1611,13 +1748,22 @@
}
out << "* " << name;
}
- } else if (type->Is<sem::Sampler>()) {
+ return true;
+ }
+
+ if (type->Is<sem::Sampler>()) {
out << "sampler";
- } else if (auto* str = type->As<sem::Struct>()) {
+ return true;
+ }
+
+ if (auto* str = type->As<sem::Struct>()) {
// The struct type emits as just the name. The declaration would be emitted
// as part of emitting the declared types.
out << program_->Symbols().NameFor(str->Declaration()->name());
- } else if (auto* tex = type->As<sem::Texture>()) {
+ return true;
+ }
+
+ if (auto* tex = type->As<sem::Texture>()) {
if (tex->Is<sem::DepthTexture>()) {
out << "depth";
} else {
@@ -1684,23 +1830,30 @@
return false;
}
out << ">";
+ return true;
+ }
- } else if (type->Is<sem::U32>()) {
+ if (type->Is<sem::U32>()) {
out << "uint";
- } else if (auto* vec = type->As<sem::Vector>()) {
+ return true;
+ }
+
+ if (auto* vec = type->As<sem::Vector>()) {
if (!EmitType(out, vec->type(), "")) {
return false;
}
out << vec->size();
- } else if (type->Is<sem::Void>()) {
- out << "void";
- } else {
- diagnostics_.add_error(diag::System::Writer,
- "unknown type in EmitType: " + type->type_name());
- return false;
+ return true;
}
- return true;
+ if (type->Is<sem::Void>()) {
+ out << "void";
+ return true;
+ }
+
+ diagnostics_.add_error(diag::System::Writer,
+ "unknown type in EmitType: " + type->type_name());
+ return false;
}
bool GeneratorImpl::EmitPackedType(std::ostream& out,
@@ -2039,6 +2192,10 @@
return SizeAndAlign{str->Size(), str->Align()};
}
+ if (auto* atomic = ty->As<sem::Atomic>()) {
+ return MslPackedTypeSizeAndAlign(atomic->Type());
+ }
+
TINT_UNREACHABLE(Writer, diagnostics_)
<< "Unhandled type " << ty->TypeInfo().name;
return {};
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index a527eb9..0f5bf1a 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -104,6 +104,15 @@
bool EmitIntrinsicCall(std::ostream& out,
ast::CallExpression* expr,
const sem::Intrinsic* intrinsic);
+ /// Handles generating a call to an atomic function (`atomicAdd`,
+ /// `atomicMax`, etc)
+ /// @param out the output of the expression stream
+ /// @param expr the call expression
+ /// @param intrinsic the semantic information for the atomic intrinsic
+ /// @returns true if the call expression is emitted
+ bool EmitAtomicCall(std::ostream& out,
+ ast::CallExpression* expr,
+ const sem::Intrinsic* intrinsic);
/// Handles generating a call to a texture function (`textureSample`,
/// `textureSampleGrad`, etc)
/// @param out the output of the expression stream
@@ -263,24 +272,6 @@
ast::InterpolationSampling sampling) const;
private:
- /// @returns the resolved type of the ast::Expression `expr`
- /// @param expr the expression
- sem::Type* TypeOf(ast::Expression* expr) const {
- return program_->TypeOf(expr);
- }
-
- /// @returns the resolved type of the ast::Type `type`
- /// @param type the type
- const sem::Type* TypeOf(const ast::Type* type) const {
- return program_->TypeOf(type);
- }
-
- /// @returns the resolved type of the ast::TypeDecl `type_decl`
- /// @param type_decl the type declaration
- const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
- return program_->TypeOf(type_decl);
- }
-
// A pair of byte size and alignment `uint32_t`s.
struct SizeAndAlign {
uint32_t size;
@@ -291,7 +282,6 @@
/// type.
SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
- const Program* program_ = nullptr;
std::function<bool()> emit_continuing_;
};
diff --git a/src/writer/text_generator.cc b/src/writer/text_generator.cc
index a3ec05c..6a6c364 100644
--- a/src/writer/text_generator.cc
+++ b/src/writer/text_generator.cc
@@ -14,10 +14,13 @@
#include "src/writer/text_generator.h"
+#include <limits>
+
namespace tint {
namespace writer {
-TextGenerator::TextGenerator() = default;
+TextGenerator::TextGenerator(const Program* program)
+ : program_(program), builder_(ProgramBuilder::Wrap(program)) {}
TextGenerator::~TextGenerator() = default;
@@ -31,6 +34,10 @@
}
}
+std::string TextGenerator::UniqueIdentifier(const std::string& prefix) {
+ return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
+}
+
TextGenerator::LineWriter::LineWriter(TextGenerator* generator)
: gen(generator) {}
diff --git a/src/writer/text_generator.h b/src/writer/text_generator.h
index 8f68385..57011d9 100644
--- a/src/writer/text_generator.h
+++ b/src/writer/text_generator.h
@@ -20,6 +20,7 @@
#include <utility>
#include "src/diagnostic/diagnostic.h"
+#include "src/program_builder.h"
namespace tint {
namespace writer {
@@ -28,7 +29,8 @@
class TextGenerator {
public:
/// Constructor
- TextGenerator();
+ /// @param program the program used by the generator
+ explicit TextGenerator(const Program* program);
~TextGenerator();
/// Increment the emitter indent level
@@ -58,6 +60,11 @@
/// @returns the error
std::string error() const { return diagnostics_.str(); }
+ /// @return a new, unique identifier with the given prefix.
+ /// @param prefix optional prefix to apply to the generated identifier. If
+ /// empty "tint" will be used.
+ std::string UniqueIdentifier(const std::string& prefix = "");
+
protected:
/// LineWriter is a helper that acts as a string buffer, who's content is
/// emitted to the TextGenerator as a single line on destruction.
@@ -122,9 +129,31 @@
TextGenerator* gen;
};
+ /// @returns the resolved type of the ast::Expression `expr`
+ /// @param expr the expression
+ sem::Type* TypeOf(ast::Expression* expr) const {
+ return builder_.TypeOf(expr);
+ }
+
+ /// @returns the resolved type of the ast::Type `type`
+ /// @param type the type
+ const sem::Type* TypeOf(const ast::Type* type) const {
+ return builder_.TypeOf(type);
+ }
+
+ /// @returns the resolved type of the ast::TypeDecl `type_decl`
+ /// @param type_decl the type
+ const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
+ return builder_.TypeOf(type_decl);
+ }
+
/// @returns a new LineWriter, used for buffering and writing a line to out_
LineWriter line() { return LineWriter(this); }
+ /// The program
+ Program const* const program_;
+ /// A ProgramBuilder that thinly wraps program_
+ ProgramBuilder builder_;
/// The text output stream
std::ostringstream out_;
/// Diagnostics generated by the generator
diff --git a/src/writer/text_generator_test.cc b/src/writer/text_generator_test.cc
new file mode 100644
index 0000000..10c7086
--- /dev/null
+++ b/src/writer/text_generator_test.cc
@@ -0,0 +1,48 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/writer/text_generator.h"
+
+#include "gtest/gtest.h"
+
+namespace tint {
+namespace writer {
+namespace {
+
+TEST(TextGeneratorTest, UniqueIdentifier) {
+ Program program(ProgramBuilder{});
+
+ TextGenerator gen(&program);
+
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident");
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_1");
+}
+
+TEST(TextGeneratorTest, UniqueIdentifier_ConflictWithExisting) {
+ ProgramBuilder builder;
+ builder.Symbols().Register("ident_1");
+ builder.Symbols().Register("ident_2");
+ Program program(std::move(builder));
+
+ TextGenerator gen(&program);
+
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident");
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_3");
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_4");
+ ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_5");
+}
+
+} // namespace
+} // namespace writer
+} // namespace tint
diff --git a/src/writer/wgsl/generator.cc b/src/writer/wgsl/generator.cc
index f190463..36eeabc 100644
--- a/src/writer/wgsl/generator.cc
+++ b/src/writer/wgsl/generator.cc
@@ -13,6 +13,7 @@
// limitations under the License.
#include "src/writer/wgsl/generator.h"
+#include "src/writer/wgsl/generator_impl.h"
namespace tint {
namespace writer {
diff --git a/src/writer/wgsl/generator.h b/src/writer/wgsl/generator.h
index 6085ecf..6b0b358 100644
--- a/src/writer/wgsl/generator.h
+++ b/src/writer/wgsl/generator.h
@@ -19,12 +19,13 @@
#include <string>
#include "src/writer/text.h"
-#include "src/writer/wgsl/generator_impl.h"
namespace tint {
namespace writer {
namespace wgsl {
+class GeneratorImpl;
+
/// Class to generate WGSL source
class Generator : public Text {
public:
@@ -46,6 +47,9 @@
std::string error() const;
private:
+ Generator(const Generator&) = delete;
+ Generator& operator=(const Generator&) = delete;
+
std::unique_ptr<GeneratorImpl> impl_;
};
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 42785a8..077e990 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -15,7 +15,6 @@
#include "src/writer/wgsl/generator_impl.h"
#include <algorithm>
-#include <limits>
#include "src/ast/access.h"
#include "src/ast/alias.h"
@@ -60,8 +59,7 @@
namespace writer {
namespace wgsl {
-GeneratorImpl::GeneratorImpl(const Program* program)
- : TextGenerator(), program_(program) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
GeneratorImpl::~GeneratorImpl() = default;
@@ -1059,24 +1057,6 @@
return true;
}
-std::string GeneratorImpl::UniqueIdentifier(const std::string& suffix) {
- auto const limit =
- std::numeric_limits<decltype(next_unique_identifier_suffix)>::max();
- while (next_unique_identifier_suffix < limit) {
- auto ident = "tint_" + std::to_string(next_unique_identifier_suffix);
- if (!suffix.empty()) {
- ident += "_" + suffix;
- }
- next_unique_identifier_suffix++;
- if (!program_->Symbols().Get(ident).IsValid()) {
- return ident;
- }
- }
- diagnostics_.add_error(diag::System::Writer,
- "Unable to generate a unique WGSL identifier");
- return "<invalid-ident>";
-}
-
} // namespace wgsl
} // namespace writer
} // namespace tint
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h
index ad9f535..7d7c732 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -194,13 +194,6 @@
/// @param decos the decoration list
/// @returns true if the decorations were emitted
bool EmitDecorations(const ast::DecorationList& decos);
-
- private:
- /// @return a new, unique, valid WGSL identifier with the given suffix.
- std::string UniqueIdentifier(const std::string& suffix = "");
-
- Program const* const program_;
- uint32_t next_unique_identifier_suffix = 0;
};
} // namespace wgsl
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index 991dc69..3273e92 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/src/writer/wgsl/generator_impl_type_test.cc
@@ -141,10 +141,10 @@
ASSERT_TRUE(gen.EmitStructType(s)) << gen.error();
EXPECT_EQ(gen.result(), R"(struct S {
[[size(8)]]
- tint_0_padding : u32;
+ padding : u32;
a : i32;
[[size(4)]]
- tint_1_padding : u32;
+ padding_1 : u32;
b : f32;
};
)");
@@ -162,10 +162,10 @@
ASSERT_TRUE(gen.EmitStructType(s)) << gen.error();
EXPECT_EQ(gen.result(), R"(struct S {
[[size(8)]]
- tint_1_padding : u32;
+ padding : u32;
tint_0_padding : i32;
[[size(4)]]
- tint_3_padding : u32;
+ padding_1 : u32;
tint_2_padding : f32;
};
)");
diff --git a/test/BUILD.gn b/test/BUILD.gn
index e120d8b..fba8220 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -305,6 +305,7 @@
"../src/utils/unique_vector_test.cc",
"../src/writer/append_vector_test.cc",
"../src/writer/float_to_string_test.cc",
+ "../src/writer/text_generator_test.cc",
]
deps = [
diff --git a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
index 9cf781e..825608b 100644
--- a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicAdd_794055(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicAdd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAdd_794055(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicAdd_794055(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicAdd_794055(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
index 86b5492..4d1bb52 100644
--- a/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_8a199a() {
- var res : u32 = atomicAdd(&(sb_rw.arg_0), 1u);
+void atomicAdd_8a199a(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicAdd_8a199a();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAdd_8a199a(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicAdd_8a199a();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAdd_8a199a(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
index 85b0ee5..a1e838a 100644
--- a/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_d32fe4() {
- var res : i32 = atomicAdd(&(sb_rw.arg_0), 1);
+void atomicAdd_d32fe4(device SB_RW& sb_rw) {
+ int res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicAdd_d32fe4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAdd_d32fe4(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicAdd_d32fe4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAdd_d32fe4(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
index 3e13cfd..ee51545 100644
--- a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicAdd_d5db1d(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicAdd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAdd_d5db1d(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicAdd_d5db1d(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicAdd_d5db1d(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
index efcd326..d095ebb 100644
--- a/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_152966() {
- var res : i32 = atomicAnd(&(sb_rw.arg_0), 1);
+void atomicAnd_152966(device SB_RW& sb_rw) {
+ int res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicAnd_152966();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAnd_152966(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicAnd_152966();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAnd_152966(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
index 6a5bb56..dd88dd3 100644
--- a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicAnd_34edd3(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicAnd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAnd_34edd3(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicAnd_34edd3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicAnd_34edd3(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
index 47755f6..7a700e9 100644
--- a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicAnd_45a819(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicAnd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAnd_45a819(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicAnd_45a819(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicAnd_45a819(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
index 31f25e8..b9ca1b0 100644
--- a/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_85a8d9() {
- var res : u32 = atomicAnd(&(sb_rw.arg_0), 1u);
+void atomicAnd_85a8d9(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicAnd_85a8d9();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAnd_85a8d9(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicAnd_85a8d9();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicAnd_85a8d9(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
index 7135dea..56d8090 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_12871c() {
- var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+void atomicCompareExchangeWeak_12871c(device SB_RW& sb_rw) {
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+ int2 res = int2(prev_value, matched);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicCompareExchangeWeak_12871c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicCompareExchangeWeak_12871c(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicCompareExchangeWeak_12871c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicCompareExchangeWeak_12871c(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
index a6c99ce..c4999ee 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_6673da() {
- var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+void atomicCompareExchangeWeak_6673da(device SB_RW& sb_rw) {
+ uint prev_value = 1u;
+ bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+ uint2 res = uint2(prev_value, matched);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicCompareExchangeWeak_6673da();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicCompareExchangeWeak_6673da(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicCompareExchangeWeak_6673da();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicCompareExchangeWeak_6673da(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
index c6431fb..036e9bf 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicCompareExchangeWeak_89ea3b(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : vec2<i32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1, 1);
+using namespace metal;
+void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol_1) {
+ int prev_value = 1;
+ bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+ int2 res = int2(prev_value, matched);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicCompareExchangeWeak_89ea3b(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicCompareExchangeWeak_89ea3b(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
index 1b632df..45d921c 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicCompareExchangeWeak_b2ab2c(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : vec2<u32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1u, 1u);
+using namespace metal;
+void atomicCompareExchangeWeak_b2ab2c(threadgroup atomic_uint* const tint_symbol_1) {
+ uint prev_value = 1u;
+ bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+ uint2 res = uint2(prev_value, matched);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
index 070beca..1dafa27 100644
--- a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicExchange_0a5dca(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicExchange(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicExchange_0a5dca(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicExchange_0a5dca(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicExchange_0a5dca(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
index d3d01f5..248b62a 100644
--- a/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_d59712() {
- var res : u32 = atomicExchange(&(sb_rw.arg_0), 1u);
+void atomicExchange_d59712(device SB_RW& sb_rw) {
+ uint res = atomic_exchange_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicExchange_d59712();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicExchange_d59712(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicExchange_d59712();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicExchange_d59712(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
index 9e1b24d..5104612 100644
--- a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicExchange_e114ba(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicExchange(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicExchange_e114ba(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicExchange_e114ba(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicExchange_e114ba(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
index be0f897..9e0dd6b 100644
--- a/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_f2e22f() {
- var res : i32 = atomicExchange(&(sb_rw.arg_0), 1);
+void atomicExchange_f2e22f(device SB_RW& sb_rw) {
+ int res = atomic_exchange_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicExchange_f2e22f();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicExchange_f2e22f(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicExchange_f2e22f();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicExchange_f2e22f(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
index 39eff93..8cc7837 100644
--- a/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_0806ad() {
- var res : i32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_0806ad(device SB_RW& sb_rw) {
+ int res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicLoad_0806ad();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicLoad_0806ad(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicLoad_0806ad();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicLoad_0806ad(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
index 7d5661a..b7c08b5 100644
--- a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicLoad_361bf1(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_361bf1(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicLoad_361bf1(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicLoad_361bf1(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
index 6304ed0..7b6310a 100644
--- a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicLoad_afcc03(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_afcc03(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicLoad_afcc03(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicLoad_afcc03(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
index 2c1eeff..aef98d3 100644
--- a/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_fe6cc3() {
- var res : u32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_fe6cc3(device SB_RW& sb_rw) {
+ uint res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicLoad_fe6cc3();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicLoad_fe6cc3(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicLoad_fe6cc3();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicLoad_fe6cc3(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
index 9a9fb1d..8d45595 100644
--- a/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_51b9be() {
- var res : u32 = atomicMax(&(sb_rw.arg_0), 1u);
+void atomicMax_51b9be(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicMax_51b9be();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMax_51b9be(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicMax_51b9be();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMax_51b9be(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
index a2997a1..12babbd 100644
--- a/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_92aa72() {
- var res : i32 = atomicMax(&(sb_rw.arg_0), 1);
+void atomicMax_92aa72(device SB_RW& sb_rw) {
+ int res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicMax_92aa72();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMax_92aa72(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicMax_92aa72();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMax_92aa72(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
index bdb46ff..26d0ee6 100644
--- a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicMax_a89cc3(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicMax(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMax_a89cc3(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicMax_a89cc3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicMax_a89cc3(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
index 5dc602c..dd139ec 100644
--- a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicMax_beccfc(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicMax(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMax_beccfc(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicMax_beccfc(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicMax_beccfc(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
index 7eb65c5..282c957 100644
--- a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicMin_278235(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicMin(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMin_278235(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicMin_278235(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicMin_278235(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
index 5ed9a71..4e6ec92 100644
--- a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicMin_69d383(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicMin(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMin_69d383(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicMin_69d383(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicMin_69d383(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
index dc15ddd..5de8f10 100644
--- a/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_8e38dc() {
- var res : i32 = atomicMin(&(sb_rw.arg_0), 1);
+void atomicMin_8e38dc(device SB_RW& sb_rw) {
+ int res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicMin_8e38dc();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMin_8e38dc(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicMin_8e38dc();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMin_8e38dc(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
index 1b28fe8..04d9d54 100644
--- a/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_c67a74() {
- var res : u32 = atomicMin(&(sb_rw.arg_0), 1u);
+void atomicMin_c67a74(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicMin_c67a74();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMin_c67a74(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicMin_c67a74();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicMin_c67a74(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
index e38ef78..3f85fc0 100644
--- a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicOr_5e3d61(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicOr(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicOr_5e3d61(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicOr_5e3d61(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicOr_5e3d61(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
index e8ae23b..a711f2d 100644
--- a/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_5e95d4() {
- var res : u32 = atomicOr(&(sb_rw.arg_0), 1u);
+void atomicOr_5e95d4(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicOr_5e95d4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicOr_5e95d4(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicOr_5e95d4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicOr_5e95d4(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
index 22c839d..3602800 100644
--- a/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_8d96a0() {
- var res : i32 = atomicOr(&(sb_rw.arg_0), 1);
+void atomicOr_8d96a0(device SB_RW& sb_rw) {
+ int res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicOr_8d96a0();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicOr_8d96a0(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicOr_8d96a0();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicOr_8d96a0(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
index e8c75dd..620f392 100644
--- a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicOr_d09248(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicOr(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicOr_d09248(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicOr_d09248(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicOr_d09248(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
index c2468d6..bebd8e5 100644
--- a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicStore_726882(tint_symbol : ptr<workgroup, atomic<u32>>) {
- atomicStore(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicStore_726882(threadgroup atomic_uint* const tint_symbol_1) {
+ atomic_store_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicStore_726882(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicStore_726882(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
index 651a1da..9f9249a 100644
--- a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicStore_8bea94(tint_symbol : ptr<workgroup, atomic<i32>>) {
- atomicStore(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicStore_8bea94(threadgroup atomic_int* const tint_symbol_1) {
+ atomic_store_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicStore_8bea94(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicStore_8bea94(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
index cefa61e..24231c1 100644
--- a/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_cdc29e() {
- atomicStore(&(sb_rw.arg_0), 1u);
+void atomicStore_cdc29e(device SB_RW& sb_rw) {
+ atomic_store_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicStore_cdc29e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicStore_cdc29e(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicStore_cdc29e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicStore_cdc29e(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
index 98cae7a..4d3c468 100644
--- a/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_d1e9a6() {
- atomicStore(&(sb_rw.arg_0), 1);
+void atomicStore_d1e9a6(device SB_RW& sb_rw) {
+ atomic_store_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicStore_d1e9a6();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicStore_d1e9a6(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicStore_d1e9a6();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicStore_d1e9a6(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
index 9bccb9b..c43fc6b 100644
--- a/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<u32>;
+ /* 0x0000 */ atomic_uint arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_54510e() {
- var res : u32 = atomicXor(&(sb_rw.arg_0), 1u);
+void atomicXor_54510e(device SB_RW& sb_rw) {
+ uint res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicXor_54510e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicXor_54510e(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicXor_54510e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicXor_54510e(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
index 32bbb93..e1b8265 100644
--- a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicXor_75dc95(tint_symbol : ptr<workgroup, atomic<i32>>) {
- var res : i32 = atomicXor(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicXor_75dc95(threadgroup atomic_int* const tint_symbol_1) {
+ int res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
- atomicXor_75dc95(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_int tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicXor_75dc95(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
index bd565c8..d48ee31 100644
--- a/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-[[block]]
+using namespace metal;
struct SB_RW {
- arg_0 : atomic<i32>;
+ /* 0x0000 */ atomic_int arg_0;
};
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_c1b78c() {
- var res : i32 = atomicXor(&(sb_rw.arg_0), 1);
+void atomicXor_c1b78c(device SB_RW& sb_rw) {
+ int res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
}
-[[stage(fragment)]]
-fn fragment_main() {
- atomicXor_c1b78c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicXor_c1b78c(sb_rw);
+ return;
}
-[[stage(compute)]]
-fn compute_main() {
- atomicXor_c1b78c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+ atomicXor_c1b78c(sb_rw);
+ return;
}
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
index 4ee0172..3f0283d 100644
--- a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
-
-fn atomicXor_c8e6be(tint_symbol : ptr<workgroup, atomic<u32>>) {
- var res : u32 = atomicXor(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicXor_c8e6be(threadgroup atomic_uint* const tint_symbol_1) {
+ uint res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
}
-[[stage(compute)]]
-fn compute_main() {
- [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
- atomicXor_c8e6be(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup atomic_uint tint_symbol_2;
+ if ((local_invocation_index == 0u)) {
+ atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomicXor_c8e6be(&(tint_symbol_2));
+ return;
}
-error: cannot declare an atomic var in a function scope
