tint: Add support for atomic ops to spirv reader

The following operations are supported:

OpAtomicLoad
OpAtomicStore
OpAtomicExchange
OpAtomicCompareExchange
OpAtomicCompareExchangeWeak
OpAtomicIIncrement
OpAtomicIDecrement
OpAtomicIAdd
OpAtomicISub
OpAtomicSMin
OpAtomicUMin
OpAtomicSMax
OpAtomicUMax
OpAtomicAnd
OpAtomicOr
OpAtomicXor

These are not, but may be supported in the future:

OpAtomicFlagTestAndSet
OpAtomicFlagClear
OpAtomicFMinEXT
OpAtomicFMaxEXT
OpAtomicFAddEXT

Bug: tint:1441
Change-Id: Ifd53643b38d43664905a0dddfca609add4914670
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94121
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

src/tint/BUILD.gn

@@ -526,6 +526,8 @@ libtint_source_set("libtint_core_all_src") {
     "transform/simplify_pointers.h",
     "transform/single_entry_point.cc",
     "transform/single_entry_point.h",
+    "transform/spirv_atomic.cc",
+    "transform/spirv_atomic.h",
     "transform/transform.cc",
     "transform/transform.h",
     "transform/unshadow.cc",
@@ -1190,6 +1192,7 @@ if (tint_build_unittests) {
       "transform/robustness_test.cc",
       "transform/simplify_pointers_test.cc",
       "transform/single_entry_point_test.cc",
+      "transform/spirv_atomic_test.cc",
       "transform/test_helper.h",
       "transform/transform_test.cc",
       "transform/unshadow_test.cc",

src/tint/CMakeLists.txt

@@ -450,6 +450,8 @@ set(TINT_LIB_SRCS
   transform/simplify_pointers.h
   transform/single_entry_point.cc
   transform/single_entry_point.h
+  transform/spirv_atomic.cc
+  transform/spirv_atomic.h
   transform/transform.cc
   transform/transform.h
   transform/unshadow.cc
@@ -1114,6 +1116,7 @@ if(TINT_BUILD_TESTS)
       transform/robustness_test.cc
       transform/simplify_pointers_test.cc
       transform/single_entry_point_test.cc
+      transform/spirv_atomic_test.cc
       transform/test_helper.h
       transform/unshadow_test.cc
       transform/unwind_discard_functions_test.cc

src/tint/reader/spirv/function.cc

@@ -36,6 +36,7 @@
 #include "src/tint/sem/builtin_type.h"
 #include "src/tint/sem/depth_texture.h"
 #include "src/tint/sem/sampled_texture.h"
+#include "src/tint/transform/spirv_atomic.h"
 
 // Terms:
 //    CFG: the control flow graph of the function, where basic blocks are the
@@ -500,6 +501,38 @@ bool IsSampledImageAccess(SpvOp opcode) {
     return false;
 }
 
+// @param opcode a SPIR-V opcode
+// @returns true if the given instruction is an atomic operation.
+bool IsAtomicOp(SpvOp opcode) {
+    switch (opcode) {
+        case SpvOpAtomicLoad:
+        case SpvOpAtomicStore:
+        case SpvOpAtomicExchange:
+        case SpvOpAtomicCompareExchange:
+        case SpvOpAtomicCompareExchangeWeak:
+        case SpvOpAtomicIIncrement:
+        case SpvOpAtomicIDecrement:
+        case SpvOpAtomicIAdd:
+        case SpvOpAtomicISub:
+        case SpvOpAtomicSMin:
+        case SpvOpAtomicUMin:
+        case SpvOpAtomicSMax:
+        case SpvOpAtomicUMax:
+        case SpvOpAtomicAnd:
+        case SpvOpAtomicOr:
+        case SpvOpAtomicXor:
+        case SpvOpAtomicFlagTestAndSet:
+        case SpvOpAtomicFlagClear:
+        case SpvOpAtomicFMinEXT:
+        case SpvOpAtomicFMaxEXT:
+        case SpvOpAtomicFAddEXT:
+            return true;
+        default:
+            break;
+    }
+    return false;
+}
+
 // @param opcode a SPIR-V opcode
 // @returns true if the given instruction is an image sampling, gather,
 // or gather-compare operation.
@@ -3487,6 +3520,10 @@ bool FunctionEmitter::EmitStatement(const spvtools::opt::Instruction& inst) {
         return EmitImageAccess(inst);
     }
 
+    if (IsAtomicOp(inst.opcode())) {
+        return EmitAtomicOp(inst);
+    }
+
     switch (inst.opcode()) {
         case SpvOpNop:
             return true;
@@ -5417,6 +5454,115 @@ bool FunctionEmitter::EmitImageQuery(const spvtools::opt::Instruction& inst) {
     return Fail() << "unhandled image query: " << inst.PrettyPrint();
 }
 
+bool FunctionEmitter::EmitAtomicOp(const spvtools::opt::Instruction& inst) {
+    auto emit_atomic = [&](sem::BuiltinType builtin, std::initializer_list<TypedExpression> args) {
+        // Split args into params and expressions
+        ast::ParameterList params;
+        params.reserve(args.size());
+        ast::ExpressionList exprs;
+        exprs.reserve(args.size());
+        size_t i = 0;
+        for (auto& a : args) {
+            params.emplace_back(builder_.Param("p" + std::to_string(i++), a.type->Build(builder_)));
+            exprs.emplace_back(a.expr);
+        }
+
+        // Function return type
+        const ast::Type* ret_type = nullptr;
+        if (inst.type_id() != 0) {
+            ret_type = parser_impl_.ConvertType(inst.type_id())->Build(builder_);
+        } else {
+            ret_type = builder_.ty.void_();
+        }
+
+        // Emit stub, will be removed by transform::SpirvAtomic
+        auto sym = builder_.Symbols().New(std::string("stub_") + sem::str(builtin));
+        auto* stub_deco =
+            builder_.ASTNodes().Create<transform::SpirvAtomic::Stub>(builder_.ID(), builtin);
+        auto* stub =
+            create<ast::Function>(Source{}, sym, std::move(params), ret_type,
+                                  /* body */ nullptr,
+                                  ast::AttributeList{
+                                      stub_deco,
+                                      builder_.Disable(ast::DisabledValidation::kFunctionHasNoBody),
+                                  },
+                                  ast::AttributeList{});
+        builder_.AST().AddFunction(stub);
+
+        // Emit call to stub, will be replaced with call to atomic builtin by transform::SpirvAtomic
+        auto* call = builder_.Call(Source{}, sym, exprs);
+        if (inst.type_id() != 0) {
+            auto* result_type = parser_impl_.ConvertType(inst.type_id());
+            TypedExpression expr{result_type, call};
+            return EmitConstDefOrWriteToHoistedVar(inst, expr);
+        }
+        AddStatement(create<ast::CallStatement>(call));
+
+        return true;
+    };
+
+    auto oper = [&](uint32_t index) -> TypedExpression {  //
+        return MakeOperand(inst, index);
+    };
+
+    auto lit = [&](int v) -> TypedExpression {
+        auto* result_type = parser_impl_.ConvertType(inst.type_id());
+        if (result_type->Is<I32>()) {
+            return TypedExpression(result_type, builder_.Expr(i32(v)));
+        } else if (result_type->Is<U32>()) {
+            return TypedExpression(result_type, builder_.Expr(u32(v)));
+        }
+        return {};
+    };
+
+    switch (inst.opcode()) {
+        case SpvOpAtomicLoad:
+            return emit_atomic(sem::BuiltinType::kAtomicLoad, {oper(/*ptr*/ 0)});
+        case SpvOpAtomicStore:
+            return emit_atomic(sem::BuiltinType::kAtomicStore,
+                               {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicExchange:
+            return emit_atomic(sem::BuiltinType::kAtomicExchange,
+                               {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicCompareExchange:
+        case SpvOpAtomicCompareExchangeWeak:
+            return emit_atomic(sem::BuiltinType::kAtomicCompareExchangeWeak,
+                               {oper(/*ptr*/ 0), /*value*/ oper(5), /*comparator*/ oper(4)});
+        case SpvOpAtomicIIncrement:
+            return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), lit(1)});
+        case SpvOpAtomicIDecrement:
+            return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), lit(1)});
+        case SpvOpAtomicIAdd:
+            return emit_atomic(sem::BuiltinType::kAtomicAdd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicISub:
+            return emit_atomic(sem::BuiltinType::kAtomicSub, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicSMin:
+            return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicUMin:
+            return emit_atomic(sem::BuiltinType::kAtomicMin, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicSMax:
+            return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicUMax:
+            return emit_atomic(sem::BuiltinType::kAtomicMax, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicAnd:
+            return emit_atomic(sem::BuiltinType::kAtomicAnd, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicOr:
+            return emit_atomic(sem::BuiltinType::kAtomicOr, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicXor:
+            return emit_atomic(sem::BuiltinType::kAtomicXor, {oper(/*ptr*/ 0), oper(/*value*/ 3)});
+        case SpvOpAtomicFlagTestAndSet:
+        case SpvOpAtomicFlagClear:
+        case SpvOpAtomicFMinEXT:
+        case SpvOpAtomicFMaxEXT:
+        case SpvOpAtomicFAddEXT:
+            return Fail() << "unsupported atomic op: " << inst.PrettyPrint();
+
+        default:
+            break;
+    }
+    return Fail() << "unhandled atomic op: " << inst.PrettyPrint();
+}
+
 ast::ExpressionList FunctionEmitter::MakeCoordinateOperandsForImageAccess(
     const spvtools::opt::Instruction& inst) {
     if (!parser_impl_.success()) {

src/tint/reader/spirv/function.h

@@ -1060,7 +1060,7 @@ class FunctionEmitter {
     /// Emits a texture builtin function call for a SPIR-V instruction that
     /// accesses an image or sampled image.
     /// @param inst the SPIR-V instruction
-    /// @returns an expression
+    /// @returns true on success, false on error
     bool EmitImageAccess(const spvtools::opt::Instruction& inst);
 
     /// Emits statements to implement a SPIR-V image query.
@@ -1068,6 +1068,11 @@ class FunctionEmitter {
     /// @returns an expression
     bool EmitImageQuery(const spvtools::opt::Instruction& inst);
 
+    /// Emits statements to implement a SPIR-V atomic op.
+    /// @param inst the SPIR-V instruction
+    /// @returns true on success, false on error
+    bool EmitAtomicOp(const spvtools::opt::Instruction& inst);
+
     /// Converts the given texel to match the type required for the storage
     /// texture with the given type. In WGSL the texel value is always provided
     /// as a 4-element vector, but the component type is determined by the

src/tint/reader/spirv/parser.cc

@@ -22,6 +22,7 @@
 #include "src/tint/transform/manager.h"
 #include "src/tint/transform/remove_unreachable_statements.h"
 #include "src/tint/transform/simplify_pointers.h"
+#include "src/tint/transform/spirv_atomic.h"
 #include "src/tint/transform/unshadow.h"
 
 namespace tint::reader::spirv {
@@ -55,6 +56,7 @@ Program Parse(const std::vector<uint32_t>& input) {
     manager.Add<transform::DecomposeStridedMatrix>();
     manager.Add<transform::DecomposeStridedArray>();
     manager.Add<transform::RemoveUnreachableStatements>();
+    manager.Add<transform::SpirvAtomic>();
     return manager.Run(&program).program;
 }
 

src/tint/tint.natvis

@@ -101,11 +101,7 @@
 		<DisplayString>{*variable};</DisplayString>
 	</Type>
 
-	<Type Name="tint::ast::UintLiteralExpression">
+	<Type Name="tint::ast::IntLiteralExpression">
-		<DisplayString>{value}</DisplayString>
-	</Type>
-
-	<Type Name="tint::ast::SintLiteralExpression">
 		<DisplayString>{value}</DisplayString>
 	</Type>
 

src/tint/transform/spirv_atomic.cc

@@ -0,0 +1,231 @@
+// Copyright 2022 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/tint/transform/spirv_atomic.h"
+
+#include <string>
+#include <unordered_map>
+#include <unordered_set>
+#include <utility>
+#include <vector>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/sem/block_statement.h"
+#include "src/tint/sem/function.h"
+#include "src/tint/sem/index_accessor_expression.h"
+#include "src/tint/sem/member_accessor_expression.h"
+#include "src/tint/sem/reference.h"
+#include "src/tint/sem/statement.h"
+#include "src/tint/utils/map.h"
+#include "src/tint/utils/unique_vector.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::SpirvAtomic);
+TINT_INSTANTIATE_TYPEINFO(tint::transform::SpirvAtomic::Stub);
+
+namespace tint::transform {
+
+/// Private implementation of transform
+struct SpirvAtomic::State {
+  private:
+    /// A struct that has been forked because a subset of members were made atomic.
+    struct ForkedStruct {
+        Symbol name;
+        std::unordered_set<size_t> atomic_members;
+    };
+
+    CloneContext& ctx;
+    ProgramBuilder& b = *ctx.dst;
+    std::unordered_map<const ast::Struct*, ForkedStruct> forked_structs;
+    std::unordered_set<const sem::Variable*> atomic_variables;
+    utils::UniqueVector<const sem::Expression*> atomic_expressions;
+
+  public:
+    /// Constructor
+    /// @param c the clone context
+    explicit State(CloneContext& c) : ctx(c) {}
+
+    /// Runs the transform
+    void Run() {
+        // Look for stub functions generated by the SPIR-V reader, which are used as placeholders
+        // for atomic builtin calls.
+        for (auto* fn : ctx.src->AST().Functions()) {
+            if (auto* stub = ast::GetAttribute<Stub>(fn->attributes)) {
+                auto* sem = ctx.src->Sem().Get(fn);
+
+                for (auto* call : sem->CallSites()) {
+                    // The first argument is always the atomic.
+                    // The stub passes this by value, whereas the builtin wants a pointer.
+                    // Take the address of the atomic argument.
+                    auto& args = call->Declaration()->args;
+                    auto out_args = ctx.Clone(args);
+                    out_args[0] = b.AddressOf(out_args[0]);
+
+                    // Replace all callsites of this stub to a call to the real builtin
+                    if (stub->builtin == sem::BuiltinType::kAtomicCompareExchangeWeak) {
+                        // atomicCompareExchangeWeak returns a struct, so insert a call to it above
+                        // the current statement, and replace the current call with the struct's
+                        // `old_value` member.
+                        auto* block = call->Stmt()->Block()->Declaration();
+                        auto old_value = b.Symbols().New("old_value");
+                        auto old_value_decl = b.Decl(b.Let(
+                            old_value, nullptr,
+                            b.MemberAccessor(b.Call(sem::str(stub->builtin), std::move(out_args)),
+                                             b.Expr("old_value"))));
+                        ctx.InsertBefore(block->statements, call->Stmt()->Declaration(),
+                                         old_value_decl);
+                        ctx.Replace(call->Declaration(), b.Expr(old_value));
+                    } else {
+                        ctx.Replace(call->Declaration(),
+                                    b.Call(sem::str(stub->builtin), std::move(out_args)));
+                    }
+
+                    // Keep track of this expression. We'll need to modify the source variable /
+                    // structure to be atomic.
+                    atomic_expressions.add(ctx.src->Sem().Get(args[0]));
+                }
+
+                // Remove the stub from the output program
+                ctx.Remove(ctx.src->AST().GlobalDeclarations(), fn);
+            }
+        }
+
+        // Transform all variables and structure members that were used in atomic operations as
+        // atomic types. This propagates up originating expression chains.
+        ProcessAtomicExpressions();
+
+        // If we need to change structure members, then fork them.
+        if (!forked_structs.empty()) {
+            ctx.ReplaceAll([&](const ast::Struct* str) {
+                // Always emit the original structure. This allows unrelated usage of the structure
+                // to continue working.
+                // auto* original = ctx.CloneWithoutTransform(str);
+
+                // Is `str` a structure we need to fork?
+                if (auto it = forked_structs.find(str); it != forked_structs.end()) {
+                    const auto& forked = it->second;
+
+                    // Re-create the structure swapping in the atomic-flavoured members
+                    std::vector<const ast::StructMember*> members(str->members.size());
+                    for (size_t i = 0; i < str->members.size(); i++) {
+                        auto* member = str->members[i];
+                        if (forked.atomic_members.count(i)) {
+                            auto* type = AtomicTypeFor(ctx.src->Sem().Get(member)->Type());
+                            auto name = ctx.src->Symbols().NameFor(member->symbol);
+                            members[i] = b.Member(name, type, ctx.Clone(member->attributes));
+                        } else {
+                            members[i] = ctx.Clone(member);
+                        }
+                    }
+                    b.Structure(forked.name, std::move(members));
+                }
+
+                // return original;
+                return nullptr;
+            });
+        }
+
+        ctx.Clone();
+    }
+
+  private:
+    ForkedStruct& Fork(const ast::Struct* str) {
+        auto& forked = forked_structs[str];
+        if (!forked.name.IsValid()) {
+            forked.name = b.Symbols().New(ctx.src->Symbols().NameFor(str->name) + "_atomic");
+        }
+        return forked;
+    }
+
+    void ProcessAtomicExpressions() {
+        for (size_t i = 0; i < atomic_expressions.size(); i++) {
+            Switch(
+                atomic_expressions[i],  //
+                [&](const sem::VariableUser* user) {
+                    auto* v = user->Variable()->Declaration();
+                    if (v->type && atomic_variables.emplace(user->Variable()).second) {
+                        ctx.Replace(v->type, AtomicTypeFor(user->Variable()->Type()));
+                    }
+                    if (auto* ctor = user->Variable()->Constructor()) {
+                        atomic_expressions.add(ctor);
+                    }
+                },
+                [&](const sem::StructMemberAccess* access) {
+                    // Fork the struct (the first time) and mark member(s) that need to be made
+                    // atomic.
+                    auto* member = access->Member();
+                    Fork(member->Struct()->Declaration()).atomic_members.emplace(member->Index());
+                    atomic_expressions.add(access->Object());
+                },
+                [&](const sem::IndexAccessorExpression* index) {
+                    atomic_expressions.add(index->Object());
+                },
+                [&](const sem::Expression* e) {
+                    if (auto* unary = e->Declaration()->As<ast::UnaryOpExpression>()) {
+                        atomic_expressions.add(ctx.src->Sem().Get(unary->expr));
+                    }
+                });
+        }
+    }
+
+    const ast::Type* AtomicTypeFor(const sem::Type* ty) {
+        return Switch(
+            ty,  //
+            [&](const sem::I32*) { return b.ty.atomic(CreateASTTypeFor(ctx, ty)); },
+            [&](const sem::U32*) { return b.ty.atomic(CreateASTTypeFor(ctx, ty)); },
+            [&](const sem::Struct* str) { return b.ty.type_name(Fork(str->Declaration()).name); },
+            [&](const sem::Array* arr) {
+                return arr->IsRuntimeSized()
+                           ? b.ty.array(AtomicTypeFor(arr->ElemType()))
+                           : b.ty.array(AtomicTypeFor(arr->ElemType()), u32(arr->Count()));
+            },
+            [&](const sem::Pointer* ptr) {
+                return b.ty.pointer(AtomicTypeFor(ptr->StoreType()), ptr->StorageClass(),
+                                    ptr->Access());
+            },
+            [&](const sem::Reference* ref) { return AtomicTypeFor(ref->StoreType()); },
+            [&](Default) {
+                TINT_ICE(Transform, b.Diagnostics())
+                    << "unhandled type: " << ty->FriendlyName(ctx.src->Symbols());
+                return nullptr;
+            });
+    }
+};
+
+SpirvAtomic::SpirvAtomic() = default;
+SpirvAtomic::~SpirvAtomic() = default;
+
+SpirvAtomic::Stub::Stub(ProgramID pid, sem::BuiltinType b) : Base(pid), builtin(b) {}
+SpirvAtomic::Stub::~Stub() = default;
+std::string SpirvAtomic::Stub::InternalName() const {
+    return "@internal(spirv-atomic " + std::string(sem::str(builtin)) + ")";
+}
+
+const SpirvAtomic::Stub* SpirvAtomic::Stub::Clone(CloneContext* ctx) const {
+    return ctx->dst->ASTNodes().Create<SpirvAtomic::Stub>(ctx->dst->ID(), builtin);
+}
+
+bool SpirvAtomic::ShouldRun(const Program* program, const DataMap&) const {
+    for (auto* fn : program->AST().Functions()) {
+        if (ast::HasAttribute<Stub>(fn->attributes)) {
+            return true;
+        }
+    }
+    return false;
+}
+
+void SpirvAtomic::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
+    State{ctx}.Run();
+}
+
+}  // namespace tint::transform

src/tint/transform/spirv_atomic.h

@@ -0,0 +1,84 @@
+// Copyright 2022 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.
+
+#ifndef SRC_TINT_TRANSFORM_SPIRV_ATOMIC_H_
+#define SRC_TINT_TRANSFORM_SPIRV_ATOMIC_H_
+
+#include <string>
+
+#include "src/tint/ast/internal_attribute.h"
+#include "src/tint/sem/builtin_type.h"
+#include "src/tint/transform/transform.h"
+
+// Forward declarations
+namespace tint {
+class CloneContext;
+}  // namespace tint
+
+namespace tint::transform {
+
+/// SpirvAtomic is a transform that replaces calls to stub functions created by the SPIR-V reader
+/// with calls to the WGSL atomic builtin. It also makes sure to replace variable declarations that
+/// are the target of the atomic operations with an atomic declaration of the same type. For
+/// structs, it creates a copy of the original struct with atomic members.
+class SpirvAtomic final : public Castable<SpirvAtomic, Transform> {
+  public:
+    /// Constructor
+    SpirvAtomic();
+    /// Destructor
+    ~SpirvAtomic() override;
+
+    /// Stub is an attribute applied to stub SPIR-V reader generated functions that need to be
+    /// translated to an atomic builtin.
+    class Stub final : public Castable<Stub, ast::InternalAttribute> {
+      public:
+        /// @param program_id the identifier of the program that owns this node
+        /// @param builtin the atomic builtin this stub represents
+        Stub(ProgramID program_id, sem::BuiltinType builtin);
+        /// Destructor
+        ~Stub() override;
+
+        /// @return a short description of the internal attribute which will be
+        /// displayed as `@internal(<name>)`
+        std::string InternalName() const override;
+
+        /// Performs a deep clone of this object using the CloneContext `ctx`.
+        /// @param ctx the clone context
+        /// @return the newly cloned object
+        const Stub* Clone(CloneContext* ctx) const override;
+
+        /// The type of the intrinsic
+        const sem::BuiltinType builtin;
+    };
+
+    /// @param program the program to inspect
+    /// @param data optional extra transform-specific input data
+    /// @returns true if this transform should be run for the given program
+    bool ShouldRun(const Program* program, const DataMap& data = {}) const override;
+
+  protected:
+    struct State;
+
+    /// Runs the transform using the CloneContext built for transforming a
+    /// program. Run() is responsible for calling Clone() on the CloneContext.
+    /// @param ctx the CloneContext primed with the input program and
+    /// ProgramBuilder
+    /// @param inputs optional extra transform-specific input data
+    /// @param outputs optional extra transform-specific output data
+    void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override;
+};
+
+}  // namespace tint::transform
+
+#endif  // SRC_TINT_TRANSFORM_SPIRV_ATOMIC_H_

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %19 = OpConstantNull %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %51 = OpConstantNull %int
+         %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+         %16 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %19
+               OpStore %idx %local_invocation_index
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %21 %22 None
+               OpBranch %23
+         %23 = OpLabel
+         %25 = OpLoad %uint %idx
+         %27 = OpULessThan %bool %25 %uint_6
+         %24 = OpLogicalNot %bool %27
+               OpSelectionMerge %29 None
+               OpBranchConditional %24 %30 %29
+         %30 = OpLabel
+               OpBranch %21
+         %29 = OpLabel
+         %31 = OpLoad %uint %idx
+         %32 = OpUDiv %uint %31 %uint_2
+         %33 = OpLoad %uint %idx
+         %34 = OpUMod %uint %33 %uint_2
+         %35 = OpLoad %uint %idx
+         %36 = OpUMod %uint %35 %uint_1
+         %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+               OpAtomicStore %41 %uint_2 %uint_0 %19
+               OpBranch %22
+         %22 = OpLabel
+         %42 = OpLoad %uint %idx
+         %43 = OpIAdd %uint %42 %uint_1
+               OpStore %idx %43
+               OpBranch %20
+         %21 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+               OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %53
+         %55 = OpLabel
+         %57 = OpLoad %uint %local_invocation_index_1
+         %56 = OpFunctionCall %void %compute_main_inner %57
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.glsl

@@ -0,0 +1,41 @@
+#version 310 es
+
+uint local_invocation_index_1 = 0u;
+shared uint wg[3][2][1];
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    for(; !(!((idx < 6u))); idx = (idx + 1u)) {
+      atomicExchange(wg[(idx / 2u)][(idx % 2u)][(idx % 1u)], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[2][1][0], 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+      uint i = (idx_1 / 2u);
+      uint i_1 = (idx_1 % 2u);
+      uint i_2 = (idx_1 % 1u);
+      atomicExchange(wg[i][i_1][i_2], 0u);
+    }
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.hlsl

@@ -0,0 +1,47 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    [loop] for(; !(!((idx < 6u))); idx = (idx + 1u)) {
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[(idx / 2u)][(idx % 2u)][(idx % 1u)], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    [loop] for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+      const uint i = (idx_1 / 2u);
+      const uint i_1 = (idx_1 % 2u);
+      const uint i_2 = (idx_1 % 1u);
+      uint atomic_result_2 = 0u;
+      InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result_2);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl

@@ -0,0 +1,74 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  uint arr[1];
+};
+
+struct tint_array_wrapper_1 {
+  tint_array_wrapper arr[2];
+};
+
+struct tint_array_wrapper_2 {
+  tint_array_wrapper_1 arr[3];
+};
+
+struct tint_array_wrapper_5 {
+  atomic_uint arr[1];
+};
+
+struct tint_array_wrapper_4 {
+  tint_array_wrapper_5 arr[2];
+};
+
+struct tint_array_wrapper_3 {
+  tint_array_wrapper_4 arr[3];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  while (true) {
+    uint const x_25 = idx;
+    if (!((x_25 < 6u))) {
+      break;
+    }
+    uint const x_31 = idx;
+    uint const x_33 = idx;
+    uint const x_35 = idx;
+    atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+    {
+      uint const x_42 = idx;
+      idx = (x_42 + 1u);
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+  uint const x_57 = *(tint_symbol_1);
+  compute_main_inner(x_57, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+    uint const i = (idx_1 / 2u);
+    uint const i_1 = (idx_1 % 2u);
+    uint const i_2 = (idx_1 % 1u);
+    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper_3 tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.spvasm

@@ -0,0 +1,140 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 90
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %idx_1 "idx_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %53 = OpConstantNull %int
+         %55 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %15
+%local_invocation_index = OpFunctionParameter %uint
+         %19 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx %6
+               OpStore %idx %local_invocation_index
+               OpBranch %22
+         %22 = OpLabel
+               OpLoopMerge %23 %24 None
+               OpBranch %25
+         %25 = OpLabel
+         %26 = OpLoad %uint %idx
+         %29 = OpULessThan %bool %26 %uint_6
+         %27 = OpLogicalNot %bool %29
+               OpSelectionMerge %31 None
+               OpBranchConditional %27 %32 %31
+         %32 = OpLabel
+               OpBranch %23
+         %31 = OpLabel
+         %33 = OpLoad %uint %idx
+         %34 = OpLoad %uint %idx
+         %35 = OpLoad %uint %idx
+         %39 = OpUDiv %uint %33 %uint_2
+         %40 = OpUMod %uint %34 %uint_2
+         %41 = OpUMod %uint %35 %uint_1
+         %43 = OpAccessChain %_ptr_Workgroup_uint %wg %39 %40 %41
+               OpAtomicStore %43 %uint_2 %uint_0 %6
+               OpBranch %24
+         %24 = OpLabel
+         %44 = OpLoad %uint %idx
+         %45 = OpIAdd %uint %44 %uint_1
+               OpStore %idx %45
+               OpBranch %22
+         %23 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %54 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %53
+               OpAtomicStore %54 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %55
+         %57 = OpLabel
+         %58 = OpLoad %uint %local_invocation_index_1
+         %59 = OpFunctionCall %void %compute_main_inner %58
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %15
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %62 = OpLabel
+      %idx_1 = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx_1 %local_invocation_index_1_param
+               OpBranch %64
+         %64 = OpLabel
+               OpLoopMerge %65 %66 None
+               OpBranch %67
+         %67 = OpLabel
+         %69 = OpLoad %uint %idx_1
+         %70 = OpULessThan %bool %69 %uint_6
+         %68 = OpLogicalNot %bool %70
+               OpSelectionMerge %71 None
+               OpBranchConditional %68 %72 %71
+         %72 = OpLabel
+               OpBranch %65
+         %71 = OpLabel
+         %73 = OpLoad %uint %idx_1
+         %74 = OpUDiv %uint %73 %uint_2
+         %75 = OpLoad %uint %idx_1
+         %76 = OpUMod %uint %75 %uint_2
+         %77 = OpLoad %uint %idx_1
+         %78 = OpUMod %uint %77 %uint_1
+         %81 = OpAccessChain %_ptr_Workgroup_uint %wg %74 %76 %78
+               OpAtomicStore %81 %uint_2 %uint_0 %6
+               OpBranch %66
+         %66 = OpLabel
+         %82 = OpLoad %uint %idx_1
+         %83 = OpIAdd %uint %82 %uint_1
+               OpStore %idx_1 %83
+               OpBranch %64
+         %65 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %85 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %55
+         %87 = OpLabel
+         %89 = OpLoad %uint %local_invocation_index_1_param_1
+         %88 = OpFunctionCall %void %compute_main_inner_1 %89
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.wgsl

@@ -0,0 +1,44 @@
+type Arr = array<u32, 1u>;
+
+type Arr_1 = array<Arr, 2u>;
+
+type Arr_2 = array<Arr_1, 3u>;
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : array<array<array<atomic<u32>, 1u>, 2u>, 3u>;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  var idx : u32 = 0u;
+  idx = local_invocation_index;
+  loop {
+    let x_25 : u32 = idx;
+    if (!((x_25 < 6u))) {
+      break;
+    }
+    let x_31 : u32 = idx;
+    let x_33 : u32 = idx;
+    let x_35 : u32 = idx;
+    atomicStore(&(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u);
+
+    continuing {
+      let x_42 : u32 = idx;
+      idx = (x_42 + 1u);
+    }
+  }
+  workgroupBarrier();
+  atomicStore(&(wg[2i][1i][0i]), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_57 : u32 = local_invocation_index_1;
+  compute_main_inner(x_57);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/array/aliased_arrays.wgsl

@@ -0,0 +1,11 @@
+type A0 = atomic<u32>;
+type A1 = array<A0, 1>;
+type A2 = array<A1, 2>;
+type A3 = array<A2, 3>;
+
+var<workgroup> wg : A3;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg[2][1][0], 1u);
+}

test/tint/builtins/atomicStore/array/array.spvasm

@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_4 ArrayStride 4
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
+         %wg = OpVariable %_ptr_Workgroup__arr_uint_uint_4 Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %15 = OpConstantNull %uint
+       %bool = OpTypeBool
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+     %uint_1 = OpConstant %uint 1
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+         %43 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %15
+               OpStore %idx %local_invocation_index
+               OpBranch %16
+         %16 = OpLabel
+               OpLoopMerge %17 %18 None
+               OpBranch %19
+         %19 = OpLabel
+         %21 = OpLoad %uint %idx
+         %22 = OpULessThan %bool %21 %uint_4
+         %20 = OpLogicalNot %bool %22
+               OpSelectionMerge %24 None
+               OpBranchConditional %20 %25 %24
+         %25 = OpLabel
+               OpBranch %17
+         %24 = OpLabel
+         %26 = OpLoad %uint %idx
+         %32 = OpAccessChain %_ptr_Workgroup_uint %wg %26
+               OpAtomicStore %32 %uint_2 %uint_0 %15
+               OpBranch %18
+         %18 = OpLabel
+         %33 = OpLoad %uint %idx
+         %35 = OpIAdd %uint %33 %uint_1
+               OpStore %idx %35
+               OpBranch %16
+         %17 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %42 = OpAccessChain %_ptr_Workgroup_uint %wg %int_1
+               OpAtomicStore %42 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %43
+         %45 = OpLabel
+         %47 = OpLoad %uint %local_invocation_index_1
+         %46 = OpFunctionCall %void %compute_main_inner %47
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/array.spvasm.expected.glsl

@@ -0,0 +1,39 @@
+#version 310 es
+
+uint local_invocation_index_1 = 0u;
+shared uint wg[4];
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    for(; !(!((idx < 4u))); idx = (idx + 1u)) {
+      atomicExchange(wg[idx], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[1], 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
+      uint i = idx_1;
+      atomicExchange(wg[i], 0u);
+    }
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/array/array.spvasm.expected.hlsl

@@ -0,0 +1,45 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[4];
+
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    [loop] for(; !(!((idx < 4u))); idx = (idx + 1u)) {
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[idx], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[1], 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    [loop] for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
+      const uint i = idx_1;
+      uint atomic_result_2 = 0u;
+      InterlockedExchange(wg[i], 0u, atomic_result_2);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/array/array.spvasm.expected.msl

@@ -0,0 +1,54 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  uint arr[4];
+};
+
+struct tint_array_wrapper_1 {
+  atomic_uint arr[4];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  while (true) {
+    uint const x_21 = idx;
+    if (!((x_21 < 4u))) {
+      break;
+    }
+    uint const x_26 = idx;
+    atomic_store_explicit(&((*(tint_symbol)).arr[x_26]), 0u, memory_order_relaxed);
+    {
+      uint const x_33 = idx;
+      idx = (x_33 + 1u);
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).arr[1]), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+  uint const x_47 = *(tint_symbol_1);
+  compute_main_inner(x_47, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
+    uint const i = idx_1;
+    atomic_store_explicit(&((*(tint_symbol_3)).arr[i]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper_1 tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/array/array.spvasm.expected.spvasm

@@ -0,0 +1,123 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 75
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %idx_1 "idx_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_4 ArrayStride 4
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+     %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
+         %wg = OpVariable %_ptr_Workgroup__arr_uint_uint_4 Workgroup
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+       %bool = OpTypeBool
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+     %uint_1 = OpConstant %uint 1
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+         %45 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %11
+%local_invocation_index = OpFunctionParameter %uint
+         %15 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx %6
+               OpStore %idx %local_invocation_index
+               OpBranch %18
+         %18 = OpLabel
+               OpLoopMerge %19 %20 None
+               OpBranch %21
+         %21 = OpLabel
+         %22 = OpLoad %uint %idx
+         %24 = OpULessThan %bool %22 %uint_4
+         %23 = OpLogicalNot %bool %24
+               OpSelectionMerge %26 None
+               OpBranchConditional %23 %27 %26
+         %27 = OpLabel
+               OpBranch %19
+         %26 = OpLabel
+         %28 = OpLoad %uint %idx
+         %34 = OpAccessChain %_ptr_Workgroup_uint %wg %28
+               OpAtomicStore %34 %uint_2 %uint_0 %6
+               OpBranch %20
+         %20 = OpLabel
+         %35 = OpLoad %uint %idx
+         %37 = OpIAdd %uint %35 %uint_1
+               OpStore %idx %37
+               OpBranch %18
+         %19 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %44 = OpAccessChain %_ptr_Workgroup_uint %wg %int_1
+               OpAtomicStore %44 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %45
+         %47 = OpLabel
+         %48 = OpLoad %uint %local_invocation_index_1
+         %49 = OpFunctionCall %void %compute_main_inner %48
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %11
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %52 = OpLabel
+      %idx_1 = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx_1 %local_invocation_index_1_param
+               OpBranch %54
+         %54 = OpLabel
+               OpLoopMerge %55 %56 None
+               OpBranch %57
+         %57 = OpLabel
+         %59 = OpLoad %uint %idx_1
+         %60 = OpULessThan %bool %59 %uint_4
+         %58 = OpLogicalNot %bool %60
+               OpSelectionMerge %61 None
+               OpBranchConditional %58 %62 %61
+         %62 = OpLabel
+               OpBranch %55
+         %61 = OpLabel
+         %63 = OpLoad %uint %idx_1
+         %66 = OpAccessChain %_ptr_Workgroup_uint %wg %63
+               OpAtomicStore %66 %uint_2 %uint_0 %6
+               OpBranch %56
+         %56 = OpLabel
+         %67 = OpLoad %uint %idx_1
+         %68 = OpIAdd %uint %67 %uint_1
+               OpStore %idx_1 %68
+               OpBranch %54
+         %55 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %70 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %45
+         %72 = OpLabel
+         %74 = OpLoad %uint %local_invocation_index_1_param_1
+         %73 = OpFunctionCall %void %compute_main_inner_1 %74
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/array.spvasm.expected.wgsl

@@ -0,0 +1,38 @@
+type Arr = array<u32, 4u>;
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : array<atomic<u32>, 4u>;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  var idx : u32 = 0u;
+  idx = local_invocation_index;
+  loop {
+    let x_21 : u32 = idx;
+    if (!((x_21 < 4u))) {
+      break;
+    }
+    let x_26 : u32 = idx;
+    atomicStore(&(wg[x_26]), 0u);
+
+    continuing {
+      let x_33 : u32 = idx;
+      idx = (x_33 + 1u);
+    }
+  }
+  workgroupBarrier();
+  atomicStore(&(wg[1i]), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_47 : u32 = local_invocation_index_1;
+  compute_main_inner(x_47);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/array/array.wgsl

@@ -0,0 +1,6 @@
+var<workgroup> wg : array<atomic<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg[1], 1u);
+}

test/tint/builtins/atomicStore/array/arrays.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %19 = OpConstantNull %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %51 = OpConstantNull %int
+         %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+         %16 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %19
+               OpStore %idx %local_invocation_index
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %21 %22 None
+               OpBranch %23
+         %23 = OpLabel
+         %25 = OpLoad %uint %idx
+         %27 = OpULessThan %bool %25 %uint_6
+         %24 = OpLogicalNot %bool %27
+               OpSelectionMerge %29 None
+               OpBranchConditional %24 %30 %29
+         %30 = OpLabel
+               OpBranch %21
+         %29 = OpLabel
+         %31 = OpLoad %uint %idx
+         %32 = OpUDiv %uint %31 %uint_2
+         %33 = OpLoad %uint %idx
+         %34 = OpUMod %uint %33 %uint_2
+         %35 = OpLoad %uint %idx
+         %36 = OpUMod %uint %35 %uint_1
+         %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+               OpAtomicStore %41 %uint_2 %uint_0 %19
+               OpBranch %22
+         %22 = OpLabel
+         %42 = OpLoad %uint %idx
+         %43 = OpIAdd %uint %42 %uint_1
+               OpStore %idx %43
+               OpBranch %20
+         %21 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+               OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %53
+         %55 = OpLabel
+         %57 = OpLoad %uint %local_invocation_index_1
+         %56 = OpFunctionCall %void %compute_main_inner %57
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.glsl

@@ -0,0 +1,41 @@
+#version 310 es
+
+uint local_invocation_index_1 = 0u;
+shared uint wg[3][2][1];
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    for(; !(!((idx < 6u))); idx = (idx + 1u)) {
+      atomicExchange(wg[(idx / 2u)][(idx % 2u)][(idx % 1u)], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[2][1][0], 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+      uint i = (idx_1 / 2u);
+      uint i_1 = (idx_1 % 2u);
+      uint i_2 = (idx_1 % 1u);
+      atomicExchange(wg[i][i_1][i_2], 0u);
+    }
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.hlsl

@@ -0,0 +1,47 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint wg[3][2][1];
+
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    [loop] for(; !(!((idx < 6u))); idx = (idx + 1u)) {
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[(idx / 2u)][(idx % 2u)][(idx % 1u)], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    [loop] for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+      const uint i = (idx_1 / 2u);
+      const uint i_1 = (idx_1 % 2u);
+      const uint i_2 = (idx_1 % 1u);
+      uint atomic_result_2 = 0u;
+      InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result_2);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl

@@ -0,0 +1,74 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  uint arr[1];
+};
+
+struct tint_array_wrapper_1 {
+  tint_array_wrapper arr[2];
+};
+
+struct tint_array_wrapper_2 {
+  tint_array_wrapper_1 arr[3];
+};
+
+struct tint_array_wrapper_5 {
+  atomic_uint arr[1];
+};
+
+struct tint_array_wrapper_4 {
+  tint_array_wrapper_5 arr[2];
+};
+
+struct tint_array_wrapper_3 {
+  tint_array_wrapper_4 arr[3];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  while (true) {
+    uint const x_25 = idx;
+    if (!((x_25 < 6u))) {
+      break;
+    }
+    uint const x_31 = idx;
+    uint const x_33 = idx;
+    uint const x_35 = idx;
+    atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+    {
+      uint const x_42 = idx;
+      idx = (x_42 + 1u);
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+  uint const x_57 = *(tint_symbol_1);
+  compute_main_inner(x_57, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+    uint const i = (idx_1 / 2u);
+    uint const i_1 = (idx_1 % 2u);
+    uint const i_2 = (idx_1 % 1u);
+    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper_3 tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.spvasm

@@ -0,0 +1,140 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 90
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %idx_1 "idx_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %53 = OpConstantNull %int
+         %55 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %15
+%local_invocation_index = OpFunctionParameter %uint
+         %19 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx %6
+               OpStore %idx %local_invocation_index
+               OpBranch %22
+         %22 = OpLabel
+               OpLoopMerge %23 %24 None
+               OpBranch %25
+         %25 = OpLabel
+         %26 = OpLoad %uint %idx
+         %29 = OpULessThan %bool %26 %uint_6
+         %27 = OpLogicalNot %bool %29
+               OpSelectionMerge %31 None
+               OpBranchConditional %27 %32 %31
+         %32 = OpLabel
+               OpBranch %23
+         %31 = OpLabel
+         %33 = OpLoad %uint %idx
+         %34 = OpLoad %uint %idx
+         %35 = OpLoad %uint %idx
+         %39 = OpUDiv %uint %33 %uint_2
+         %40 = OpUMod %uint %34 %uint_2
+         %41 = OpUMod %uint %35 %uint_1
+         %43 = OpAccessChain %_ptr_Workgroup_uint %wg %39 %40 %41
+               OpAtomicStore %43 %uint_2 %uint_0 %6
+               OpBranch %24
+         %24 = OpLabel
+         %44 = OpLoad %uint %idx
+         %45 = OpIAdd %uint %44 %uint_1
+               OpStore %idx %45
+               OpBranch %22
+         %23 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %54 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %53
+               OpAtomicStore %54 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %55
+         %57 = OpLabel
+         %58 = OpLoad %uint %local_invocation_index_1
+         %59 = OpFunctionCall %void %compute_main_inner %58
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %15
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %62 = OpLabel
+      %idx_1 = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx_1 %local_invocation_index_1_param
+               OpBranch %64
+         %64 = OpLabel
+               OpLoopMerge %65 %66 None
+               OpBranch %67
+         %67 = OpLabel
+         %69 = OpLoad %uint %idx_1
+         %70 = OpULessThan %bool %69 %uint_6
+         %68 = OpLogicalNot %bool %70
+               OpSelectionMerge %71 None
+               OpBranchConditional %68 %72 %71
+         %72 = OpLabel
+               OpBranch %65
+         %71 = OpLabel
+         %73 = OpLoad %uint %idx_1
+         %74 = OpUDiv %uint %73 %uint_2
+         %75 = OpLoad %uint %idx_1
+         %76 = OpUMod %uint %75 %uint_2
+         %77 = OpLoad %uint %idx_1
+         %78 = OpUMod %uint %77 %uint_1
+         %81 = OpAccessChain %_ptr_Workgroup_uint %wg %74 %76 %78
+               OpAtomicStore %81 %uint_2 %uint_0 %6
+               OpBranch %66
+         %66 = OpLabel
+         %82 = OpLoad %uint %idx_1
+         %83 = OpIAdd %uint %82 %uint_1
+               OpStore %idx_1 %83
+               OpBranch %64
+         %65 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %85 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %55
+         %87 = OpLabel
+         %89 = OpLoad %uint %local_invocation_index_1_param_1
+         %88 = OpFunctionCall %void %compute_main_inner_1 %89
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/array/arrays.spvasm.expected.wgsl

@@ -0,0 +1,44 @@
+type Arr = array<u32, 1u>;
+
+type Arr_1 = array<Arr, 2u>;
+
+type Arr_2 = array<Arr_1, 3u>;
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : array<array<array<atomic<u32>, 1u>, 2u>, 3u>;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  var idx : u32 = 0u;
+  idx = local_invocation_index;
+  loop {
+    let x_25 : u32 = idx;
+    if (!((x_25 < 6u))) {
+      break;
+    }
+    let x_31 : u32 = idx;
+    let x_33 : u32 = idx;
+    let x_35 : u32 = idx;
+    atomicStore(&(wg[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u);
+
+    continuing {
+      let x_42 : u32 = idx;
+      idx = (x_42 + 1u);
+    }
+  }
+  workgroupBarrier();
+  atomicStore(&(wg[2i][1i][0i]), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_57 : u32 = local_invocation_index_1;
+  compute_main_inner(x_57);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/array/arrays.wgsl

@@ -0,0 +1,6 @@
+var<workgroup> wg : array<array<array<atomic<u32>, 1>, 2>, 3>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg[2][1][0], 1u);
+}

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm

@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+               OpDecorate %_arr_S_uint_10 ArrayStride 12
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+    %uint_10 = OpConstant %uint 10
+%_arr_S_uint_10 = OpTypeArray %S %uint_10
+%_ptr_Workgroup__arr_S_uint_10 = OpTypePointer Workgroup %_arr_S_uint_10
+         %wg = OpVariable %_ptr_Workgroup__arr_S_uint_10 Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %32 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %17
+               OpStore %idx %local_invocation_index
+               OpBranch %18
+         %18 = OpLabel
+               OpLoopMerge %19 %20 None
+               OpBranch %21
+         %21 = OpLabel
+         %23 = OpLoad %uint %idx
+         %24 = OpULessThan %bool %23 %uint_10
+         %22 = OpLogicalNot %bool %24
+               OpSelectionMerge %26 None
+               OpBranchConditional %22 %27 %26
+         %27 = OpLabel
+               OpBranch %19
+         %26 = OpLabel
+         %28 = OpLoad %uint %idx
+         %31 = OpAccessChain %_ptr_Workgroup_int %wg %28 %uint_0
+               OpStore %31 %32
+         %38 = OpAccessChain %_ptr_Workgroup_uint %wg %28 %uint_1
+               OpAtomicStore %38 %uint_2 %uint_0 %17
+         %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %28 %uint_2
+               OpStore %40 %17
+               OpBranch %20
+         %20 = OpLabel
+         %41 = OpLoad %uint %idx
+         %42 = OpIAdd %uint %41 %uint_1
+               OpStore %idx %42
+               OpBranch %18
+         %19 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %48 = OpAccessChain %_ptr_Workgroup_uint %wg %int_4 %uint_1
+               OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %49
+         %51 = OpLabel
+         %53 = OpLoad %uint %local_invocation_index_1
+         %52 = OpFunctionCall %void %compute_main_inner %53
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.glsl

@@ -0,0 +1,56 @@
+#version 310 es
+
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S_atomic wg[10];
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    for(; !(!((idx < 10u))); idx = (idx + 1u)) {
+      uint x_28 = idx;
+      wg[x_28].x = 0;
+      atomicExchange(wg[x_28].a, 0u);
+      wg[x_28].y = 0u;
+    }
+  }
+  barrier();
+  atomicExchange(wg[4].a, 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+      uint i = idx_1;
+      wg[i].x = 0;
+      atomicExchange(wg[i].a, 0u);
+      wg[i].y = 0u;
+    }
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.hlsl

@@ -0,0 +1,56 @@
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S_atomic wg[10];
+
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  {
+    [loop] for(; !(!((idx < 10u))); idx = (idx + 1u)) {
+      const uint x_28 = idx;
+      wg[x_28].x = 0;
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[x_28].a, 0u, atomic_result);
+      wg[x_28].y = 0u;
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[4].a, 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    [loop] for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+      const uint i = idx_1;
+      wg[i].x = 0;
+      uint atomic_result_2 = 0u;
+      InterlockedExchange(wg[i].a, 0u, atomic_result_2);
+      wg[i].y = 0u;
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl

@@ -0,0 +1,70 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S_atomic {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+struct tint_array_wrapper {
+  S arr[10];
+};
+
+struct tint_array_wrapper_1 {
+  S_atomic arr[10];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+  uint idx = 0u;
+  idx = local_invocation_index;
+  while (true) {
+    uint const x_23 = idx;
+    if (!((x_23 < 10u))) {
+      break;
+    }
+    uint const x_28 = idx;
+    (*(tint_symbol)).arr[x_28].x = 0;
+    atomic_store_explicit(&((*(tint_symbol)).arr[x_28].a), 0u, memory_order_relaxed);
+    (*(tint_symbol)).arr[x_28].y = 0u;
+    {
+      uint const x_41 = idx;
+      idx = (x_41 + 1u);
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).arr[4].a), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+  uint const x_53 = *(tint_symbol_1);
+  compute_main_inner(x_53, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+    uint const i = idx_1;
+    (*(tint_symbol_3)).arr[i].x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].a), 0u, memory_order_relaxed);
+    (*(tint_symbol_3)).arr[i].y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper_1 tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.spvasm

@@ -0,0 +1,142 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 83
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S_atomic "S_atomic"
+               OpMemberName %S_atomic 0 "x"
+               OpMemberName %S_atomic 1 "a"
+               OpMemberName %S_atomic 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %idx_1 "idx_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S_atomic 0 Offset 0
+               OpMemberDecorate %S_atomic 1 Offset 4
+               OpMemberDecorate %S_atomic 2 Offset 8
+               OpDecorate %_arr_S_atomic_uint_10 ArrayStride 12
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+   %S_atomic = OpTypeStruct %int %uint %uint
+    %uint_10 = OpConstant %uint 10
+%_arr_S_atomic_uint_10 = OpTypeArray %S_atomic %uint_10
+%_ptr_Workgroup__arr_S_atomic_uint_10 = OpTypePointer Workgroup %_arr_S_atomic_uint_10
+         %wg = OpVariable %_ptr_Workgroup__arr_S_atomic_uint_10 Workgroup
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %34 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %51 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %13
+%local_invocation_index = OpFunctionParameter %uint
+         %17 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx %6
+               OpStore %idx %local_invocation_index
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %21 %22 None
+               OpBranch %23
+         %23 = OpLabel
+         %24 = OpLoad %uint %idx
+         %26 = OpULessThan %bool %24 %uint_10
+         %25 = OpLogicalNot %bool %26
+               OpSelectionMerge %28 None
+               OpBranchConditional %25 %29 %28
+         %29 = OpLabel
+               OpBranch %21
+         %28 = OpLabel
+         %30 = OpLoad %uint %idx
+         %33 = OpAccessChain %_ptr_Workgroup_int %wg %30 %uint_0
+               OpStore %33 %34
+         %40 = OpAccessChain %_ptr_Workgroup_uint %wg %30 %uint_1
+               OpAtomicStore %40 %uint_2 %uint_0 %6
+         %42 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %30 %uint_2
+               OpStore %42 %6
+               OpBranch %22
+         %22 = OpLabel
+         %43 = OpLoad %uint %idx
+         %44 = OpIAdd %uint %43 %uint_1
+               OpStore %idx %44
+               OpBranch %20
+         %21 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %50 = OpAccessChain %_ptr_Workgroup_uint %wg %int_4 %uint_1
+               OpAtomicStore %50 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %51
+         %53 = OpLabel
+         %54 = OpLoad %uint %local_invocation_index_1
+         %55 = OpFunctionCall %void %compute_main_inner %54
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %13
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %58 = OpLabel
+      %idx_1 = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx_1 %local_invocation_index_1_param
+               OpBranch %60
+         %60 = OpLabel
+               OpLoopMerge %61 %62 None
+               OpBranch %63
+         %63 = OpLabel
+         %65 = OpLoad %uint %idx_1
+         %66 = OpULessThan %bool %65 %uint_10
+         %64 = OpLogicalNot %bool %66
+               OpSelectionMerge %67 None
+               OpBranchConditional %64 %68 %67
+         %68 = OpLabel
+               OpBranch %61
+         %67 = OpLabel
+         %69 = OpLoad %uint %idx_1
+         %70 = OpAccessChain %_ptr_Workgroup_int %wg %69 %uint_0
+               OpStore %70 %34
+         %73 = OpAccessChain %_ptr_Workgroup_uint %wg %69 %uint_1
+               OpAtomicStore %73 %uint_2 %uint_0 %6
+         %74 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %69 %uint_2
+               OpStore %74 %6
+               OpBranch %62
+         %62 = OpLabel
+         %75 = OpLoad %uint %idx_1
+         %76 = OpIAdd %uint %75 %uint_1
+               OpStore %idx_1 %76
+               OpBranch %60
+         %61 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %78 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %51
+         %80 = OpLabel
+         %82 = OpLoad %uint %local_invocation_index_1_param_1
+         %81 = OpFunctionCall %void %compute_main_inner_1 %82
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.wgsl

@@ -0,0 +1,52 @@
+struct S_atomic {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+struct S {
+  x : i32,
+  a : u32,
+  y : u32,
+}
+
+type Arr = array<S, 10u>;
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : array<S_atomic, 10u>;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  var idx : u32 = 0u;
+  idx = local_invocation_index;
+  loop {
+    let x_23 : u32 = idx;
+    if (!((x_23 < 10u))) {
+      break;
+    }
+    let x_28 : u32 = idx;
+    wg[x_28].x = 0i;
+    atomicStore(&(wg[x_28].a), 0u);
+    wg[x_28].y = 0u;
+
+    continuing {
+      let x_41 : u32 = idx;
+      idx = (x_41 + 1u);
+    }
+  }
+  workgroupBarrier();
+  atomicStore(&(wg[4i].a), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_53 : u32 = local_invocation_index_1;
+  compute_main_inner(x_53);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/array_of_struct.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+};
+
+var<workgroup> wg : array<S, 10>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg[4].a, 1u);
+}

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm

@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "b"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+   %uint_264 = OpConstant %uint 264
+         %35 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %26 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %26 %uint_2 %uint_0 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %31 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %31 %uint_2 %uint_0 %uint_1
+         %34 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %34 %uint_2 %uint_0 %uint_2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %35
+         %37 = OpLabel
+         %39 = OpLoad %uint %local_invocation_index_1
+         %38 = OpFunctionCall %void %compute_main_inner %39
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.glsl

@@ -0,0 +1,47 @@
+#version 310 es
+
+struct S_atomic {
+  int x;
+  uint a;
+  uint b;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint b;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S_atomic wg;
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  atomicExchange(wg.a, 0u);
+  atomicExchange(wg.b, 0u);
+  barrier();
+  atomicExchange(wg.a, 1u);
+  atomicExchange(wg.b, 2u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    atomicExchange(wg.b, 0u);
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.hlsl

@@ -0,0 +1,50 @@
+struct S_atomic {
+  int x;
+  uint a;
+  uint b;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S_atomic wg;
+
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  uint atomic_result = 0u;
+  InterlockedExchange(wg.a, 0u, atomic_result);
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.b, 0u, atomic_result_1);
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_2 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_2);
+  uint atomic_result_3 = 0u;
+  InterlockedExchange(wg.b, 2u, atomic_result_3);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    uint atomic_result_4 = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result_4);
+    uint atomic_result_5 = 0u;
+    InterlockedExchange(wg.b, 0u, atomic_result_5);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.msl

@@ -0,0 +1,49 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S_atomic {
+  int x;
+  atomic_uint a;
+  atomic_uint b;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint b;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+  (*(tint_symbol)).x = 0;
+  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol)).b), 0u, memory_order_relaxed);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol)).b), 2u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup S_atomic* const tint_symbol_2) {
+  uint const x_39 = *(tint_symbol_1);
+  compute_main_inner(x_39, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S_atomic* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  {
+    (*(tint_symbol_3)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3)).a), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3)).b), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup S_atomic tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.spvasm

@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S_atomic "S_atomic"
+               OpMemberName %S_atomic 0 "x"
+               OpMemberName %S_atomic 1 "a"
+               OpMemberName %S_atomic 2 "b"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S_atomic 0 Offset 0
+               OpMemberDecorate %S_atomic 1 Offset 4
+               OpMemberDecorate %S_atomic 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+   %S_atomic = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S_atomic = OpTypePointer Workgroup %S_atomic
+         %wg = OpVariable %_ptr_Workgroup_S_atomic Workgroup
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %19 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %37 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %11
+%local_invocation_index = OpFunctionParameter %uint
+         %15 = OpLabel
+         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %18 %19
+         %25 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %25 %uint_2 %uint_0 %6
+         %28 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %28 %uint_2 %uint_0 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %33 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %33 %uint_2 %uint_0 %uint_1
+         %36 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %36 %uint_2 %uint_0 %uint_2
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %37
+         %39 = OpLabel
+         %40 = OpLoad %uint %local_invocation_index_1
+         %41 = OpFunctionCall %void %compute_main_inner %40
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %11
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %44 = OpLabel
+         %45 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %45 %19
+         %48 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %48 %uint_2 %uint_0 %6
+         %51 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %51 %uint_2 %uint_0 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %53 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %37
+         %55 = OpLabel
+         %57 = OpLoad %uint %local_invocation_index_1_param_1
+         %56 = OpFunctionCall %void %compute_main_inner_1 %57
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.spvasm.expected.wgsl

@@ -0,0 +1,37 @@
+struct S_atomic {
+  x : i32,
+  a : atomic<u32>,
+  b : atomic<u32>,
+}
+
+struct S {
+  x : i32,
+  a : u32,
+  b : u32,
+}
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : S_atomic;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  wg.x = 0i;
+  atomicStore(&(wg.a), 0u);
+  atomicStore(&(wg.b), 0u);
+  workgroupBarrier();
+  atomicStore(&(wg.a), 1u);
+  atomicStore(&(wg.b), 2u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_39 : u32 = local_invocation_index_1;
+  compute_main_inner(x_39);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl

@@ -0,0 +1,13 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  b : atomic<u32>,
+};
+
+var<workgroup> wg: S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg.a, 1u);
+  atomicStore(&wg.b, 2u);
+}

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm

@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %25 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %31
+         %33 = OpLabel
+         %35 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %35
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.glsl

@@ -0,0 +1,46 @@
+#version 310 es
+
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S_atomic wg;
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  atomicExchange(wg.a, 0u);
+  wg.y = 0u;
+  barrier();
+  atomicExchange(wg.a, 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    wg.y = 0u;
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.hlsl

@@ -0,0 +1,46 @@
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S_atomic wg;
+
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  uint atomic_result = 0u;
+  InterlockedExchange(wg.a, 0u, atomic_result);
+  wg.y = 0u;
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    uint atomic_result_2 = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result_2);
+    wg.y = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S_atomic {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+  (*(tint_symbol)).x = 0;
+  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+  (*(tint_symbol)).y = 0u;
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup S_atomic* const tint_symbol_2) {
+  uint const x_35 = *(tint_symbol_1);
+  compute_main_inner(x_35, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S_atomic* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  {
+    (*(tint_symbol_3)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3)).a), 0u, memory_order_relaxed);
+    (*(tint_symbol_3)).y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup S_atomic tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.spvasm

@@ -0,0 +1,87 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 52
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S_atomic "S_atomic"
+               OpMemberName %S_atomic 0 "x"
+               OpMemberName %S_atomic 1 "a"
+               OpMemberName %S_atomic 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S_atomic 0 Offset 0
+               OpMemberDecorate %S_atomic 1 Offset 4
+               OpMemberDecorate %S_atomic 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+   %S_atomic = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S_atomic = OpTypePointer Workgroup %S_atomic
+         %wg = OpVariable %_ptr_Workgroup_S_atomic Workgroup
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %19 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %33 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %11
+%local_invocation_index = OpFunctionParameter %uint
+         %15 = OpLabel
+         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %18 %19
+         %25 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %25 %uint_2 %uint_0 %6
+         %27 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %27 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %32 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %32 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %33
+         %35 = OpLabel
+         %36 = OpLoad %uint %local_invocation_index_1
+         %37 = OpFunctionCall %void %compute_main_inner %36
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %11
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %40 = OpLabel
+         %41 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %41 %19
+         %44 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %44 %uint_2 %uint_0 %6
+         %45 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %45 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %47 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %33
+         %49 = OpLabel
+         %51 = OpLoad %uint %local_invocation_index_1_param_1
+         %50 = OpFunctionCall %void %compute_main_inner_1 %51
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/flat_single_atomic.spvasm.expected.wgsl

@@ -0,0 +1,36 @@
+struct S_atomic {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+struct S {
+  x : i32,
+  a : u32,
+  y : u32,
+}
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : S_atomic;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  wg.x = 0i;
+  atomicStore(&(wg.a), 0u);
+  wg.y = 0u;
+  workgroupBarrier();
+  atomicStore(&(wg.a), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_35 : u32 = local_invocation_index_1;
+  compute_main_inner(x_35);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+};
+
+var<workgroup> wg: S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg.a, 1u);
+}

test/tint/builtins/atomicStore/struct/nested.spvasm

@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S2 "S2"
+               OpMemberName %S2 0 "x"
+               OpMemberName %S2 1 "y"
+               OpMemberName %S2 2 "z"
+               OpMemberName %S2 3 "a"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "x"
+               OpMemberName %S1 1 "a"
+               OpName %S0 "S0"
+               OpMemberName %S0 0 "x"
+               OpMemberName %S0 1 "a"
+               OpMemberName %S0 2 "y"
+               OpMemberName %S0 3 "z"
+               OpMemberName %S1 2 "y"
+               OpMemberName %S1 3 "z"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S2 0 Offset 0
+               OpMemberDecorate %S2 1 Offset 4
+               OpMemberDecorate %S2 2 Offset 8
+               OpMemberDecorate %S2 3 Offset 12
+               OpMemberDecorate %S1 0 Offset 0
+               OpMemberDecorate %S1 1 Offset 4
+               OpMemberDecorate %S0 0 Offset 0
+               OpMemberDecorate %S0 1 Offset 4
+               OpMemberDecorate %S0 2 Offset 8
+               OpMemberDecorate %S0 3 Offset 12
+               OpMemberDecorate %S1 2 Offset 20
+               OpMemberDecorate %S1 3 Offset 24
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+         %S0 = OpTypeStruct %int %uint %int %int
+         %S1 = OpTypeStruct %int %S0 %int %int
+         %S2 = OpTypeStruct %int %int %int %S1
+%_ptr_Workgroup_S2 = OpTypePointer Workgroup %S2
+         %wg = OpVariable %_ptr_Workgroup_S2 Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %18 = OpConstantNull %int
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %30 = OpConstantNull %uint
+   %uint_264 = OpConstant %uint 264
+         %40 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %17 %18
+         %20 = OpAccessChain %_ptr_Workgroup_int %wg %uint_1
+               OpStore %20 %18
+         %22 = OpAccessChain %_ptr_Workgroup_int %wg %uint_2
+               OpStore %22 %18
+         %24 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_0
+               OpStore %24 %18
+         %25 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_0
+               OpStore %25 %18
+         %29 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %29 %uint_2 %uint_0 %30
+         %31 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_2
+               OpStore %31 %18
+         %32 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_3
+               OpStore %32 %18
+         %33 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_2
+               OpStore %33 %18
+         %34 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_3
+               OpStore %34 %18
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %39 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %39 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %40
+         %42 = OpLabel
+         %44 = OpLoad %uint %local_invocation_index_1
+         %43 = OpFunctionCall %void %compute_main_inner %44
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/nested.spvasm.expected.glsl

@@ -0,0 +1,90 @@
+#version 310 es
+
+struct S0_atomic {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+
+struct S0 {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+
+struct S1_atomic {
+  int x;
+  S0_atomic a;
+  int y;
+  int z;
+};
+
+struct S1 {
+  int x;
+  S0 a;
+  int y;
+  int z;
+};
+
+struct S2_atomic {
+  int x;
+  int y;
+  int z;
+  S1_atomic a;
+};
+
+struct S2 {
+  int x;
+  int y;
+  int z;
+  S1 a;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S2_atomic wg;
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  wg.y = 0;
+  wg.z = 0;
+  wg.a.x = 0;
+  wg.a.a.x = 0;
+  atomicExchange(wg.a.a.a, 0u);
+  wg.a.a.y = 0;
+  wg.a.a.z = 0;
+  wg.a.y = 0;
+  wg.a.z = 0;
+  barrier();
+  atomicExchange(wg.a.a.a, 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    wg.y = 0;
+    wg.z = 0;
+    wg.a.x = 0;
+    wg.a.a.x = 0;
+    atomicExchange(wg.a.a.a, 0u);
+    wg.a.a.y = 0;
+    wg.a.a.z = 0;
+    wg.a.y = 0;
+    wg.a.z = 0;
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/nested.spvasm.expected.hlsl

@@ -0,0 +1,73 @@
+struct S0_atomic {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+struct S1_atomic {
+  int x;
+  S0_atomic a;
+  int y;
+  int z;
+};
+struct S2_atomic {
+  int x;
+  int y;
+  int z;
+  S1_atomic a;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S2_atomic wg;
+
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  wg.y = 0;
+  wg.z = 0;
+  wg.a.x = 0;
+  wg.a.a.x = 0;
+  uint atomic_result = 0u;
+  InterlockedExchange(wg.a.a.a, 0u, atomic_result);
+  wg.a.a.y = 0;
+  wg.a.a.z = 0;
+  wg.a.y = 0;
+  wg.a.z = 0;
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a.a.a, 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    wg.y = 0;
+    wg.z = 0;
+    wg.a.x = 0;
+    wg.a.a.x = 0;
+    uint atomic_result_2 = 0u;
+    InterlockedExchange(wg.a.a.a, 0u, atomic_result_2);
+    wg.a.a.y = 0;
+    wg.a.a.z = 0;
+    wg.a.y = 0;
+    wg.a.z = 0;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/nested.spvasm.expected.msl

@@ -0,0 +1,92 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S0_atomic {
+  int x;
+  atomic_uint a;
+  int y;
+  int z;
+};
+
+struct S0 {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+
+struct S1_atomic {
+  int x;
+  S0_atomic a;
+  int y;
+  int z;
+};
+
+struct S1 {
+  int x;
+  S0 a;
+  int y;
+  int z;
+};
+
+struct S2_atomic {
+  int x;
+  int y;
+  int z;
+  S1_atomic a;
+};
+
+struct S2 {
+  int x;
+  int y;
+  int z;
+  S1 a;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S2_atomic* const tint_symbol) {
+  (*(tint_symbol)).x = 0;
+  (*(tint_symbol)).y = 0;
+  (*(tint_symbol)).z = 0;
+  (*(tint_symbol)).a.x = 0;
+  (*(tint_symbol)).a.a.x = 0;
+  atomic_store_explicit(&((*(tint_symbol)).a.a.a), 0u, memory_order_relaxed);
+  (*(tint_symbol)).a.a.y = 0;
+  (*(tint_symbol)).a.a.z = 0;
+  (*(tint_symbol)).a.y = 0;
+  (*(tint_symbol)).a.z = 0;
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a.a.a), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup S2_atomic* const tint_symbol_2) {
+  uint const x_44 = *(tint_symbol_1);
+  compute_main_inner(x_44, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S2_atomic* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  {
+    (*(tint_symbol_3)).x = 0;
+    (*(tint_symbol_3)).y = 0;
+    (*(tint_symbol_3)).z = 0;
+    (*(tint_symbol_3)).a.x = 0;
+    (*(tint_symbol_3)).a.a.x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3)).a.a.a), 0u, memory_order_relaxed);
+    (*(tint_symbol_3)).a.a.y = 0;
+    (*(tint_symbol_3)).a.a.z = 0;
+    (*(tint_symbol_3)).a.y = 0;
+    (*(tint_symbol_3)).a.z = 0;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup S2_atomic tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/nested.spvasm.expected.spvasm

@@ -0,0 +1,137 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S2_atomic "S2_atomic"
+               OpMemberName %S2_atomic 0 "x"
+               OpMemberName %S2_atomic 1 "y"
+               OpMemberName %S2_atomic 2 "z"
+               OpMemberName %S2_atomic 3 "a"
+               OpName %S1_atomic "S1_atomic"
+               OpMemberName %S1_atomic 0 "x"
+               OpMemberName %S1_atomic 1 "a"
+               OpName %S0_atomic "S0_atomic"
+               OpMemberName %S0_atomic 0 "x"
+               OpMemberName %S0_atomic 1 "a"
+               OpMemberName %S0_atomic 2 "y"
+               OpMemberName %S0_atomic 3 "z"
+               OpMemberName %S1_atomic 2 "y"
+               OpMemberName %S1_atomic 3 "z"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S2_atomic 0 Offset 0
+               OpMemberDecorate %S2_atomic 1 Offset 4
+               OpMemberDecorate %S2_atomic 2 Offset 8
+               OpMemberDecorate %S2_atomic 3 Offset 12
+               OpMemberDecorate %S1_atomic 0 Offset 0
+               OpMemberDecorate %S1_atomic 1 Offset 4
+               OpMemberDecorate %S0_atomic 0 Offset 0
+               OpMemberDecorate %S0_atomic 1 Offset 4
+               OpMemberDecorate %S0_atomic 2 Offset 8
+               OpMemberDecorate %S0_atomic 3 Offset 12
+               OpMemberDecorate %S1_atomic 2 Offset 20
+               OpMemberDecorate %S1_atomic 3 Offset 24
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+  %S0_atomic = OpTypeStruct %int %uint %int %int
+  %S1_atomic = OpTypeStruct %int %S0_atomic %int %int
+  %S2_atomic = OpTypeStruct %int %int %int %S1_atomic
+%_ptr_Workgroup_S2_atomic = OpTypePointer Workgroup %S2_atomic
+         %wg = OpVariable %_ptr_Workgroup_S2_atomic Workgroup
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %21 = OpConstantNull %int
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %42 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %13
+%local_invocation_index = OpFunctionParameter %uint
+         %17 = OpLabel
+         %20 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %20 %21
+         %23 = OpAccessChain %_ptr_Workgroup_int %wg %uint_1
+               OpStore %23 %21
+         %25 = OpAccessChain %_ptr_Workgroup_int %wg %uint_2
+               OpStore %25 %21
+         %27 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_0
+               OpStore %27 %21
+         %28 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_0
+               OpStore %28 %21
+         %32 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %32 %uint_2 %uint_0 %6
+         %33 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_2
+               OpStore %33 %21
+         %34 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_3
+               OpStore %34 %21
+         %35 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_2
+               OpStore %35 %21
+         %36 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_3
+               OpStore %36 %21
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %41 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %41 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %42
+         %44 = OpLabel
+         %45 = OpLoad %uint %local_invocation_index_1
+         %46 = OpFunctionCall %void %compute_main_inner %45
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %13
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %49 = OpLabel
+         %50 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %50 %21
+         %51 = OpAccessChain %_ptr_Workgroup_int %wg %uint_1
+               OpStore %51 %21
+         %52 = OpAccessChain %_ptr_Workgroup_int %wg %uint_2
+               OpStore %52 %21
+         %53 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_0
+               OpStore %53 %21
+         %54 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_0
+               OpStore %54 %21
+         %57 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %57 %uint_2 %uint_0 %6
+         %58 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_2
+               OpStore %58 %21
+         %59 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_3
+               OpStore %59 %21
+         %60 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_2
+               OpStore %60 %21
+         %61 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_3
+               OpStore %61 %21
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %63 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %42
+         %65 = OpLabel
+         %67 = OpLoad %uint %local_invocation_index_1_param_1
+         %66 = OpFunctionCall %void %compute_main_inner_1 %67
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/nested.spvasm.expected.wgsl

@@ -0,0 +1,73 @@
+struct S0_atomic {
+  x : i32,
+  a : atomic<u32>,
+  y : i32,
+  z : i32,
+}
+
+struct S0 {
+  x : i32,
+  a : u32,
+  y : i32,
+  z : i32,
+}
+
+struct S1_atomic {
+  x : i32,
+  a : S0_atomic,
+  y : i32,
+  z : i32,
+}
+
+struct S1 {
+  x : i32,
+  a : S0,
+  y : i32,
+  z : i32,
+}
+
+struct S2_atomic {
+  x : i32,
+  y : i32,
+  z : i32,
+  a : S1_atomic,
+}
+
+struct S2 {
+  x : i32,
+  y : i32,
+  z : i32,
+  a : S1,
+}
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : S2_atomic;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  wg.x = 0i;
+  wg.y = 0i;
+  wg.z = 0i;
+  wg.a.x = 0i;
+  wg.a.a.x = 0i;
+  atomicStore(&(wg.a.a.a), 0u);
+  wg.a.a.y = 0i;
+  wg.a.a.z = 0i;
+  wg.a.y = 0i;
+  wg.a.z = 0i;
+  workgroupBarrier();
+  atomicStore(&(wg.a.a.a), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_44 : u32 = local_invocation_index_1;
+  compute_main_inner(x_44);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/nested.wgsl

@@ -0,0 +1,27 @@
+struct S0 {
+  x : i32,
+  a : atomic<u32>,
+  y : i32,
+  z : i32,
+};
+
+struct S1 {
+  x : i32,
+  a : S0,
+  y : i32,
+  z : i32,
+};
+
+struct S2 {
+  x : i32,
+  y : i32,
+  z : i32,
+  a : S1,
+};
+
+var<workgroup> wg: S2;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg.a.a.a, 1u);
+}

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm

@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpDecorate %_arr_uint_uint_10 ArrayStride 4
+               OpMemberDecorate %S 2 Offset 44
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+    %uint_10 = OpConstant %uint 10
+%_arr_uint_uint_10 = OpTypeArray %uint %uint_10
+          %S = OpTypeStruct %int %_arr_uint_uint_10 %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %18 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %22 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %22
+         %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %17 %18
+         %21 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpStore %21 %22
+               OpStore %idx %local_invocation_index
+               OpBranch %25
+         %25 = OpLabel
+               OpLoopMerge %26 %27 None
+               OpBranch %28
+         %28 = OpLabel
+         %30 = OpLoad %uint %idx
+         %31 = OpULessThan %bool %30 %uint_10
+         %29 = OpLogicalNot %bool %31
+               OpSelectionMerge %33 None
+               OpBranchConditional %29 %34 %33
+         %34 = OpLabel
+               OpBranch %26
+         %33 = OpLabel
+         %35 = OpLoad %uint %idx
+         %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %35
+               OpAtomicStore %40 %uint_2 %uint_0 %22
+               OpBranch %27
+         %27 = OpLabel
+         %41 = OpLoad %uint %idx
+         %42 = OpIAdd %uint %41 %uint_1
+               OpStore %idx %42
+               OpBranch %25
+         %26 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %48 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %int_4
+               OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %49
+         %51 = OpLabel
+         %53 = OpLoad %uint %local_invocation_index_1
+         %52 = OpFunctionCall %void %compute_main_inner %53
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.glsl

@@ -0,0 +1,57 @@
+#version 310 es
+
+struct S_atomic {
+  int x;
+  uint a[10];
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a[10];
+  uint y;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S_atomic wg;
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  wg.x = 0;
+  wg.y = 0u;
+  idx = local_invocation_index;
+  {
+    for(; !(!((idx < 10u))); idx = (idx + 1u)) {
+      atomicExchange(wg.a[idx], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg.a[4], 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    wg.y = 0u;
+  }
+  {
+    for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+      uint i = idx_1;
+      atomicExchange(wg.a[i], 0u);
+    }
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.hlsl

@@ -0,0 +1,57 @@
+struct S_atomic {
+  int x;
+  uint a[10];
+  uint y;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S_atomic wg;
+
+void compute_main_inner(uint local_invocation_index) {
+  uint idx = 0u;
+  wg.x = 0;
+  wg.y = 0u;
+  idx = local_invocation_index;
+  {
+    [loop] for(; !(!((idx < 10u))); idx = (idx + 1u)) {
+      uint atomic_result = 0u;
+      InterlockedExchange(wg.a[idx], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a[4], 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    wg.y = 0u;
+  }
+  {
+    [loop] for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+      const uint i = idx_1;
+      uint atomic_result_2 = 0u;
+      InterlockedExchange(wg.a[i], 0u, atomic_result_2);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl

@@ -0,0 +1,72 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  uint arr[10];
+};
+
+struct tint_array_wrapper_1 {
+  atomic_uint arr[10];
+};
+
+struct S_atomic {
+  int x;
+  tint_array_wrapper_1 a;
+  uint y;
+};
+
+struct S {
+  int x;
+  tint_array_wrapper a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+  uint idx = 0u;
+  (*(tint_symbol)).x = 0;
+  (*(tint_symbol)).y = 0u;
+  idx = local_invocation_index;
+  while (true) {
+    uint const x_30 = idx;
+    if (!((x_30 < 10u))) {
+      break;
+    }
+    uint const x_35 = idx;
+    atomic_store_explicit(&((*(tint_symbol)).a.arr[x_35]), 0u, memory_order_relaxed);
+    {
+      uint const x_41 = idx;
+      idx = (x_41 + 1u);
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a.arr[4]), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup S_atomic* const tint_symbol_2) {
+  uint const x_53 = *(tint_symbol_1);
+  compute_main_inner(x_53, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S_atomic* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  {
+    (*(tint_symbol_3)).x = 0;
+    (*(tint_symbol_3)).y = 0u;
+  }
+  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+    uint const i = idx_1;
+    atomic_store_explicit(&((*(tint_symbol_3)).a.arr[i]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup S_atomic tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.spvasm

@@ -0,0 +1,142 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 83
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S_atomic "S_atomic"
+               OpMemberName %S_atomic 0 "x"
+               OpMemberName %S_atomic 1 "a"
+               OpMemberName %S_atomic 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %idx_1 "idx_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S_atomic 0 Offset 0
+               OpMemberDecorate %S_atomic 1 Offset 4
+               OpDecorate %_arr_uint_uint_10 ArrayStride 4
+               OpMemberDecorate %S_atomic 2 Offset 44
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+    %uint_10 = OpConstant %uint 10
+%_arr_uint_uint_10 = OpTypeArray %uint %uint_10
+   %S_atomic = OpTypeStruct %int %_arr_uint_uint_10 %uint
+%_ptr_Workgroup_S_atomic = OpTypePointer Workgroup %S_atomic
+         %wg = OpVariable %_ptr_Workgroup_S_atomic Workgroup
+       %void = OpTypeVoid
+         %13 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %23 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %51 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %13
+%local_invocation_index = OpFunctionParameter %uint
+         %17 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %6
+               OpStore %idx %6
+         %22 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %22 %23
+         %26 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpStore %26 %6
+               OpStore %idx %local_invocation_index
+               OpBranch %27
+         %27 = OpLabel
+               OpLoopMerge %28 %29 None
+               OpBranch %30
+         %30 = OpLabel
+         %31 = OpLoad %uint %idx
+         %33 = OpULessThan %bool %31 %uint_10
+         %32 = OpLogicalNot %bool %33
+               OpSelectionMerge %35 None
+               OpBranchConditional %32 %36 %35
+         %36 = OpLabel
+               OpBranch %28
+         %35 = OpLabel
+         %37 = OpLoad %uint %idx
+         %42 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %37
+               OpAtomicStore %42 %uint_2 %uint_0 %6
+               OpBranch %29
+         %29 = OpLabel
+         %43 = OpLoad %uint %idx
+         %44 = OpIAdd %uint %43 %uint_1
+               OpStore %idx %44
+               OpBranch %27
+         %28 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %50 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %int_4
+               OpAtomicStore %50 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %51
+         %53 = OpLabel
+         %54 = OpLoad %uint %local_invocation_index_1
+         %55 = OpFunctionCall %void %compute_main_inner %54
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %13
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %58 = OpLabel
+      %idx_1 = OpVariable %_ptr_Function_uint Function %6
+         %59 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %59 %23
+         %60 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpStore %60 %6
+               OpStore %idx_1 %local_invocation_index_1_param
+               OpBranch %62
+         %62 = OpLabel
+               OpLoopMerge %63 %64 None
+               OpBranch %65
+         %65 = OpLabel
+         %67 = OpLoad %uint %idx_1
+         %68 = OpULessThan %bool %67 %uint_10
+         %66 = OpLogicalNot %bool %68
+               OpSelectionMerge %69 None
+               OpBranchConditional %66 %70 %69
+         %70 = OpLabel
+               OpBranch %63
+         %69 = OpLabel
+         %71 = OpLoad %uint %idx_1
+         %74 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %71
+               OpAtomicStore %74 %uint_2 %uint_0 %6
+               OpBranch %64
+         %64 = OpLabel
+         %75 = OpLoad %uint %idx_1
+         %76 = OpIAdd %uint %75 %uint_1
+               OpStore %idx_1 %76
+               OpBranch %62
+         %63 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %78 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %51
+         %80 = OpLabel
+         %82 = OpLoad %uint %local_invocation_index_1_param_1
+         %81 = OpFunctionCall %void %compute_main_inner_1 %82
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.wgsl

@@ -0,0 +1,52 @@
+type Arr = array<u32, 10u>;
+
+struct S_atomic {
+  x : i32,
+  a : array<atomic<u32>, 10u>,
+  y : u32,
+}
+
+struct S {
+  x : i32,
+  a : Arr,
+  y : u32,
+}
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : S_atomic;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  var idx : u32 = 0u;
+  wg.x = 0i;
+  wg.y = 0u;
+  idx = local_invocation_index;
+  loop {
+    let x_30 : u32 = idx;
+    if (!((x_30 < 10u))) {
+      break;
+    }
+    let x_35 : u32 = idx;
+    atomicStore(&(wg.a[x_35]), 0u);
+
+    continuing {
+      let x_41 : u32 = idx;
+      idx = (x_41 + 1u);
+    }
+  }
+  workgroupBarrier();
+  atomicStore(&(wg.a[4i]), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_53 : u32 = local_invocation_index_1;
+  compute_main_inner(x_53);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/struct_of_array.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : array<atomic<u32>, 10>,
+  y : u32,
+};
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&wg.a[4], 1u);
+}

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm

@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %25 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %31
+         %33 = OpLabel
+         %35 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %35
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.glsl

@@ -0,0 +1,46 @@
+#version 310 es
+
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+uint local_invocation_index_1 = 0u;
+shared S_atomic wg;
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  atomicExchange(wg.a, 0u);
+  wg.y = 0u;
+  barrier();
+  atomicExchange(wg.a, 1u);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    wg.y = 0u;
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.hlsl

@@ -0,0 +1,46 @@
+struct S_atomic {
+  int x;
+  uint a;
+  uint y;
+};
+
+static uint local_invocation_index_1 = 0u;
+groupshared S_atomic wg;
+
+void compute_main_inner(uint local_invocation_index) {
+  wg.x = 0;
+  uint atomic_result = 0u;
+  InterlockedExchange(wg.a, 0u, atomic_result);
+  wg.y = 0u;
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_1);
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    wg.x = 0;
+    uint atomic_result_2 = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result_2);
+    wg.y = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.msl

@@ -0,0 +1,48 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S_atomic {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const tint_symbol) {
+  (*(tint_symbol)).x = 0;
+  atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+  (*(tint_symbol)).y = 0u;
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup S_atomic* const tint_symbol_2) {
+  uint const x_35 = *(tint_symbol_1);
+  compute_main_inner(x_35, tint_symbol_2);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S_atomic* const tint_symbol_3, thread uint* const tint_symbol_4) {
+  {
+    (*(tint_symbol_3)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3)).a), 0u, memory_order_relaxed);
+    (*(tint_symbol_3)).y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_4) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_4, tint_symbol_3);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup S_atomic tint_symbol_5;
+  thread uint tint_symbol_6 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
+  return;
+}
+

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.spvasm

@@ -0,0 +1,87 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 52
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S_atomic "S_atomic"
+               OpMemberName %S_atomic 0 "x"
+               OpMemberName %S_atomic 1 "a"
+               OpMemberName %S_atomic 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S_atomic 0 Offset 0
+               OpMemberDecorate %S_atomic 1 Offset 4
+               OpMemberDecorate %S_atomic 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+   %S_atomic = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S_atomic = OpTypePointer Workgroup %S_atomic
+         %wg = OpVariable %_ptr_Workgroup_S_atomic Workgroup
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %19 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %33 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %11
+%local_invocation_index = OpFunctionParameter %uint
+         %15 = OpLabel
+         %18 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %18 %19
+         %25 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %25 %uint_2 %uint_0 %6
+         %27 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %27 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %32 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %32 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %33
+         %35 = OpLabel
+         %36 = OpLoad %uint %local_invocation_index_1
+         %37 = OpFunctionCall %void %compute_main_inner %36
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %11
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %40 = OpLabel
+         %41 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %41 %19
+         %44 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %44 %uint_2 %uint_0 %6
+         %45 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %45 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %47 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %33
+         %49 = OpLabel
+         %51 = OpLoad %uint %local_invocation_index_1_param_1
+         %50 = OpFunctionCall %void %compute_main_inner_1 %51
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomicStore/struct/via_ptr_let.spvasm.expected.wgsl

@@ -0,0 +1,36 @@
+struct S_atomic {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+struct S {
+  x : i32,
+  a : u32,
+  y : u32,
+}
+
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> wg : S_atomic;
+
+fn compute_main_inner(local_invocation_index : u32) {
+  wg.x = 0i;
+  atomicStore(&(wg.a), 0u);
+  wg.y = 0u;
+  workgroupBarrier();
+  atomicStore(&(wg.a), 1u);
+  return;
+}
+
+fn compute_main_1() {
+  let x_35 : u32 = local_invocation_index_1;
+  compute_main_inner(x_35);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl

@@ -0,0 +1,14 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+};
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  let p0 = &wg;
+  let p1 = &((*p0).a);
+  atomicStore(p1, 1u);
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm

@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW "SB_RW"
+               OpMemberName %SB_RW 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicAdd_d32fe4 "atomicAdd_d32fe4"
+               OpName %res "res"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW Block
+               OpMemberDecorate %SB_RW 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+        %int = OpTypeInt 32 1
+      %SB_RW = OpTypeStruct %int
+%_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %19 = OpConstantNull %int
+%atomicAdd_d32fe4 = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %19
+         %15 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+          %9 = OpAtomicIAdd %int %15 %uint_1 %uint_0 %int_1
+               OpStore %res %9
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %atomicAdd_d32fe4
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %atomicAdd_d32fe4
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm.expected.glsl

@@ -0,0 +1,68 @@
+#version 310 es
+precision mediump float;
+
+struct SB_RW_atomic {
+  int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  int arg_0;
+} sb_rw;
+void atomicAdd_d32fe4() {
+  int res = 0;
+  int x_9 = atomicAdd(sb_rw.arg_0, 1);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct SB_RW_atomic {
+  int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  int arg_0;
+} sb_rw;
+void atomicAdd_d32fe4() {
+  int res = 0;
+  int x_9 = atomicAdd(sb_rw.arg_0, 1);
+  res = x_9;
+  return;
+}
+
+void compute_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+void compute_main() {
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm.expected.hlsl

@@ -0,0 +1,36 @@
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, int value) {
+  int original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
+void atomicAdd_d32fe4() {
+  int res = 0;
+  const int x_9 = tint_atomicAdd(sb_rw, 0u, 1);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB_RW_atomic {
+  /* 0x0000 */ atomic_int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+void atomicAdd_d32fe4(device SB_RW_atomic* const tint_symbol) {
+  int res = 0;
+  int const x_9 = atomic_fetch_add_explicit(&((*(tint_symbol)).arg_0), 1, memory_order_relaxed);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1(device SB_RW_atomic* const tint_symbol_1) {
+  atomicAdd_d32fe4(tint_symbol_1);
+  return;
+}
+
+fragment void fragment_main(device SB_RW_atomic* tint_symbol_2 [[buffer(0)]]) {
+  fragment_main_1(tint_symbol_2);
+  return;
+}
+
+void compute_main_1(device SB_RW_atomic* const tint_symbol_3) {
+  atomicAdd_d32fe4(tint_symbol_3);
+  return;
+}
+
+kernel void compute_main(device SB_RW_atomic* tint_symbol_4 [[buffer(0)]]) {
+  compute_main_1(tint_symbol_4);
+  return;
+}
+

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm.expected.spvasm

@@ -0,0 +1,66 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW_atomic "SB_RW_atomic"
+               OpMemberName %SB_RW_atomic 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicAdd_d32fe4 "atomicAdd_d32fe4"
+               OpName %res "res"
+               OpName %fragment_main_1 "fragment_main_1"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW_atomic Block
+               OpMemberDecorate %SB_RW_atomic 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+        %int = OpTypeInt 32 1
+%SB_RW_atomic = OpTypeStruct %int
+%_ptr_StorageBuffer_SB_RW_atomic = OpTypePointer StorageBuffer %SB_RW_atomic
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW_atomic StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+          %9 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+%atomicAdd_d32fe4 = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %9
+               OpStore %res %9
+         %18 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+         %12 = OpAtomicIAdd %int %18 %uint_1 %uint_0 %int_1
+               OpStore %res %12
+               OpReturn
+               OpFunctionEnd
+%fragment_main_1 = OpFunction %void None %5
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %atomicAdd_d32fe4
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %fragment_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %5
+         %27 = OpLabel
+         %28 = OpFunctionCall %void %atomicAdd_d32fe4
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_i32.spvasm.expected.wgsl

@@ -0,0 +1,36 @@
+struct SB_RW_atomic {
+  arg_0 : atomic<i32>,
+}
+
+struct SB_RW {
+  arg_0 : i32,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
+
+fn atomicAdd_d32fe4() {
+  var res : i32 = 0i;
+  let x_9 : i32 = atomicAdd(&(sb_rw.arg_0), 1i);
+  res = x_9;
+  return;
+}
+
+fn fragment_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+@fragment
+fn fragment_main() {
+  fragment_main_1();
+}
+
+fn compute_main_1() {
+  atomicAdd_d32fe4();
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main() {
+  compute_main_1();
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm

@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW "SB_RW"
+               OpMemberName %SB_RW 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicAdd_8a199a "atomicAdd_8a199a"
+               OpName %res "res"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW Block
+               OpMemberDecorate %SB_RW 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+       %uint = OpTypeInt 32 0
+      %SB_RW = OpTypeStruct %uint
+%_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+%atomicAdd_8a199a = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %17
+         %14 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
+          %9 = OpAtomicIAdd %uint %14 %uint_1 %uint_0 %uint_1
+               OpStore %res %9
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %atomicAdd_8a199a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %atomicAdd_8a199a
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm.expected.glsl

@@ -0,0 +1,68 @@
+#version 310 es
+precision mediump float;
+
+struct SB_RW_atomic {
+  uint arg_0;
+};
+
+struct SB_RW {
+  uint arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  uint arg_0;
+} sb_rw;
+void atomicAdd_8a199a() {
+  uint res = 0u;
+  uint x_9 = atomicAdd(sb_rw.arg_0, 1u);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct SB_RW_atomic {
+  uint arg_0;
+};
+
+struct SB_RW {
+  uint arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  uint arg_0;
+} sb_rw;
+void atomicAdd_8a199a() {
+  uint res = 0u;
+  uint x_9 = atomicAdd(sb_rw.arg_0, 1u);
+  res = x_9;
+  return;
+}
+
+void compute_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+void compute_main() {
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm.expected.hlsl

@@ -0,0 +1,36 @@
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
+void atomicAdd_8a199a() {
+  uint res = 0u;
+  const uint x_9 = tint_atomicAdd(sb_rw, 0u, 1u);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB_RW_atomic {
+  /* 0x0000 */ atomic_uint arg_0;
+};
+
+struct SB_RW {
+  uint arg_0;
+};
+
+void atomicAdd_8a199a(device SB_RW_atomic* const tint_symbol) {
+  uint res = 0u;
+  uint const x_9 = atomic_fetch_add_explicit(&((*(tint_symbol)).arg_0), 1u, memory_order_relaxed);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1(device SB_RW_atomic* const tint_symbol_1) {
+  atomicAdd_8a199a(tint_symbol_1);
+  return;
+}
+
+fragment void fragment_main(device SB_RW_atomic* tint_symbol_2 [[buffer(0)]]) {
+  fragment_main_1(tint_symbol_2);
+  return;
+}
+
+void compute_main_1(device SB_RW_atomic* const tint_symbol_3) {
+  atomicAdd_8a199a(tint_symbol_3);
+  return;
+}
+
+kernel void compute_main(device SB_RW_atomic* tint_symbol_4 [[buffer(0)]]) {
+  compute_main_1(tint_symbol_4);
+  return;
+}
+

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW_atomic "SB_RW_atomic"
+               OpMemberName %SB_RW_atomic 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicAdd_8a199a "atomicAdd_8a199a"
+               OpName %res "res"
+               OpName %fragment_main_1 "fragment_main_1"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW_atomic Block
+               OpMemberDecorate %SB_RW_atomic 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+       %uint = OpTypeInt 32 0
+%SB_RW_atomic = OpTypeStruct %uint
+%_ptr_StorageBuffer_SB_RW_atomic = OpTypePointer StorageBuffer %SB_RW_atomic
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW_atomic StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+          %9 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%atomicAdd_8a199a = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %9
+               OpStore %res %9
+         %17 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
+         %12 = OpAtomicIAdd %uint %17 %uint_1 %uint_0 %uint_1
+               OpStore %res %12
+               OpReturn
+               OpFunctionEnd
+%fragment_main_1 = OpFunction %void None %5
+         %19 = OpLabel
+         %20 = OpFunctionCall %void %atomicAdd_8a199a
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %22 = OpLabel
+         %23 = OpFunctionCall %void %fragment_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %5
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %atomicAdd_8a199a
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/storage_u32.spvasm.expected.wgsl

@@ -0,0 +1,36 @@
+struct SB_RW_atomic {
+  arg_0 : atomic<u32>,
+}
+
+struct SB_RW {
+  arg_0 : u32,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW_atomic;
+
+fn atomicAdd_8a199a() {
+  var res : u32 = 0u;
+  let x_9 : u32 = atomicAdd(&(sb_rw.arg_0), 1u);
+  res = x_9;
+  return;
+}
+
+fn fragment_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+@fragment
+fn fragment_main() {
+  fragment_main_1();
+}
+
+fn compute_main_1() {
+  atomicAdd_8a199a();
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main() {
+  compute_main_1();
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm

@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicAdd_794055 "atomicAdd_794055"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+      %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %18 = OpConstantNull %int
+         %19 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicAdd_794055 = OpFunction %void None %7
+         %10 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %18
+         %11 = OpAtomicIAdd %int %arg_0 %uint_2 %uint_0 %int_1
+               OpStore %res %11
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %19
+%local_invocation_index = OpFunctionParameter %uint
+         %22 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %18
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %27 = OpFunctionCall %void %atomicAdd_794055
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %7
+         %29 = OpLabel
+         %31 = OpLoad %uint %local_invocation_index_1
+         %30 = OpFunctionCall %void %compute_main_inner %31
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.glsl

@@ -0,0 +1,37 @@
+#version 310 es
+
+uint local_invocation_index_1 = 0u;
+shared int arg_0;
+void atomicAdd_794055() {
+  int res = 0;
+  int x_11 = atomicAdd(arg_0, 1);
+  res = x_11;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index) {
+  atomicExchange(arg_0, 0);
+  barrier();
+  atomicAdd_794055();
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    atomicExchange(arg_0, 0);
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.hlsl

@@ -0,0 +1,44 @@
+static uint local_invocation_index_1 = 0u;
+groupshared int arg_0;
+
+void atomicAdd_794055() {
+  int res = 0;
+  int atomic_result = 0;
+  InterlockedAdd(arg_0, 1, atomic_result);
+  const int x_11 = atomic_result;
+  res = x_11;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index) {
+  int atomic_result_1 = 0;
+  InterlockedExchange(arg_0, 0, atomic_result_1);
+  GroupMemoryBarrierWithGroupSync();
+  atomicAdd_794055();
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    int atomic_result_2 = 0;
+    InterlockedExchange(arg_0, 0, atomic_result_2);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void atomicAdd_794055(threadgroup atomic_int* const tint_symbol) {
+  int res = 0;
+  int const x_11 = atomic_fetch_add_explicit(tint_symbol, 1, memory_order_relaxed);
+  res = x_11;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+  atomic_store_explicit(tint_symbol_1, 0, memory_order_relaxed);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_794055(tint_symbol_1);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_2, threadgroup atomic_int* const tint_symbol_3) {
+  uint const x_31 = *(tint_symbol_2);
+  compute_main_inner(x_31, tint_symbol_3);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup atomic_int* const tint_symbol_4, thread uint* const tint_symbol_5) {
+  {
+    atomic_store_explicit(tint_symbol_4, 0, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_5) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_5, tint_symbol_4);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_6;
+  thread uint tint_symbol_7 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_6), &(tint_symbol_7));
+  return;
+}
+

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.spvasm

@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicAdd_794055 "atomicAdd_794055"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+        %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+      %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void
+         %14 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+      %int_1 = OpConstant %int 1
+         %22 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicAdd_794055 = OpFunction %void None %10
+         %13 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %14
+               OpStore %res %14
+         %17 = OpAtomicIAdd %int %arg_0 %uint_2 %uint_0 %int_1
+               OpStore %res %17
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %22
+%local_invocation_index = OpFunctionParameter %uint
+         %25 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %14
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %30 = OpFunctionCall %void %atomicAdd_794055
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %10
+         %32 = OpLabel
+         %33 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %33
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %22
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %37 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %14
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %41 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %10
+         %43 = OpLabel
+         %45 = OpLoad %uint %local_invocation_index_1_param_1
+         %44 = OpFunctionCall %void %compute_main_inner_1 %45
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_i32.spvasm.expected.wgsl

@@ -0,0 +1,29 @@
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> arg_0 : atomic<i32>;
+
+fn atomicAdd_794055() {
+  var res : i32 = 0i;
+  let x_11 : i32 = atomicAdd(&(arg_0), 1i);
+  res = x_11;
+  return;
+}
+
+fn compute_main_inner(local_invocation_index : u32) {
+  atomicStore(&(arg_0), 0i);
+  workgroupBarrier();
+  atomicAdd_794055();
+  return;
+}
+
+fn compute_main_1() {
+  let x_31 : u32 = local_invocation_index_1;
+  compute_main_inner(x_31);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm

@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicAdd_d5db1d "atomicAdd_d5db1d"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+      %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+         %18 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicAdd_d5db1d = OpFunction %void None %6
+          %9 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %17
+         %10 = OpAtomicIAdd %uint %arg_0 %uint_2 %uint_0 %uint_1
+               OpStore %res %10
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %18
+%local_invocation_index = OpFunctionParameter %uint
+         %21 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %17
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %26 = OpFunctionCall %void %atomicAdd_d5db1d
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %6
+         %28 = OpLabel
+         %30 = OpLoad %uint %local_invocation_index_1
+         %29 = OpFunctionCall %void %compute_main_inner %30
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.glsl

@@ -0,0 +1,37 @@
+#version 310 es
+
+uint local_invocation_index_1 = 0u;
+shared uint arg_0;
+void atomicAdd_d5db1d() {
+  uint res = 0u;
+  uint x_10 = atomicAdd(arg_0, 1u);
+  res = x_10;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index) {
+  atomicExchange(arg_0, 0u);
+  barrier();
+  atomicAdd_d5db1d();
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+void compute_main(uint local_invocation_index_1_param) {
+  {
+    atomicExchange(arg_0, 0u);
+  }
+  barrier();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.hlsl

@@ -0,0 +1,44 @@
+static uint local_invocation_index_1 = 0u;
+groupshared uint arg_0;
+
+void atomicAdd_d5db1d() {
+  uint res = 0u;
+  uint atomic_result = 0u;
+  InterlockedAdd(arg_0, 1u, atomic_result);
+  const uint x_10 = atomic_result;
+  res = x_10;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index) {
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(arg_0, 0u, atomic_result_1);
+  GroupMemoryBarrierWithGroupSync();
+  atomicAdd_d5db1d();
+  return;
+}
+
+void compute_main_1() {
+  compute_main_inner(local_invocation_index_1);
+  return;
+}
+
+struct tint_symbol_1 {
+  uint local_invocation_index_1_param : SV_GroupIndex;
+};
+
+void compute_main_inner_1(uint local_invocation_index_1_param) {
+  {
+    uint atomic_result_2 = 0u;
+    InterlockedExchange(arg_0, 0u, atomic_result_2);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner_1(tint_symbol.local_invocation_index_1_param);
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.msl

@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void atomicAdd_d5db1d(threadgroup atomic_uint* const tint_symbol) {
+  uint res = 0u;
+  uint const x_10 = atomic_fetch_add_explicit(tint_symbol, 1u, memory_order_relaxed);
+  res = x_10;
+  return;
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
+  atomic_store_explicit(tint_symbol_1, 0u, memory_order_relaxed);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_d5db1d(tint_symbol_1);
+  return;
+}
+
+void compute_main_1(thread uint* const tint_symbol_2, threadgroup atomic_uint* const tint_symbol_3) {
+  uint const x_30 = *(tint_symbol_2);
+  compute_main_inner(x_30, tint_symbol_3);
+  return;
+}
+
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup atomic_uint* const tint_symbol_4, thread uint* const tint_symbol_5) {
+  {
+    atomic_store_explicit(tint_symbol_4, 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  *(tint_symbol_5) = local_invocation_index_1_param;
+  compute_main_1(tint_symbol_5, tint_symbol_4);
+}
+
+kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_6;
+  thread uint tint_symbol_7 = 0u;
+  compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_6), &(tint_symbol_7));
+  return;
+}
+

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.spvasm

@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1_param_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1_param_1 "local_invocation_index_1_param_1"
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %arg_0 "arg_0"
+               OpName %atomicAdd_d5db1d "atomicAdd_d5db1d"
+               OpName %res "res"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main_1 "compute_main_1"
+               OpName %compute_main_inner_1 "compute_main_inner_1"
+               OpName %local_invocation_index_1_param "local_invocation_index_1_param"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %6 = OpConstantNull %uint
+%local_invocation_index_1 = OpVariable %_ptr_Private_uint Private %6
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+      %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+%_ptr_Function_uint = OpTypePointer Function %uint
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+         %20 = OpTypeFunction %void %uint
+   %uint_264 = OpConstant %uint 264
+%atomicAdd_d5db1d = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function_uint Function %6
+               OpStore %res %6
+         %15 = OpAtomicIAdd %uint %arg_0 %uint_2 %uint_0 %uint_1
+               OpStore %res %15
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner = OpFunction %void None %20
+%local_invocation_index = OpFunctionParameter %uint
+         %23 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %28 = OpFunctionCall %void %atomicAdd_d5db1d
+               OpReturn
+               OpFunctionEnd
+%compute_main_1 = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpLoad %uint %local_invocation_index_1
+         %32 = OpFunctionCall %void %compute_main_inner %31
+               OpReturn
+               OpFunctionEnd
+%compute_main_inner_1 = OpFunction %void None %20
+%local_invocation_index_1_param = OpFunctionParameter %uint
+         %35 = OpLabel
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %6
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %local_invocation_index_1 %local_invocation_index_1_param
+         %39 = OpFunctionCall %void %compute_main_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %41 = OpLabel
+         %43 = OpLoad %uint %local_invocation_index_1_param_1
+         %42 = OpFunctionCall %void %compute_main_inner_1 %43
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAdd/workgroup_u32.spvasm.expected.wgsl

@@ -0,0 +1,29 @@
+var<private> local_invocation_index_1 : u32;
+
+var<workgroup> arg_0 : atomic<u32>;
+
+fn atomicAdd_d5db1d() {
+  var res : u32 = 0u;
+  let x_10 : u32 = atomicAdd(&(arg_0), 1u);
+  res = x_10;
+  return;
+}
+
+fn compute_main_inner(local_invocation_index : u32) {
+  atomicStore(&(arg_0), 0u);
+  workgroupBarrier();
+  atomicAdd_d5db1d();
+  return;
+}
+
+fn compute_main_1() {
+  let x_30 : u32 = local_invocation_index_1;
+  compute_main_inner(x_30);
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn compute_main(@builtin(local_invocation_index) local_invocation_index_1_param : u32) {
+  local_invocation_index_1 = local_invocation_index_1_param;
+  compute_main_1();
+}

test/tint/builtins/atomics/from_gen/literal/atomicAnd/storage_i32.spvasm

@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %SB_RW "SB_RW"
+               OpMemberName %SB_RW 0 "arg_0"
+               OpName %sb_rw "sb_rw"
+               OpName %atomicAnd_152966 "atomicAnd_152966"
+               OpName %res "res"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %SB_RW Block
+               OpMemberDecorate %SB_RW 0 Offset 0
+               OpDecorate %sb_rw DescriptorSet 0
+               OpDecorate %sb_rw Binding 0
+        %int = OpTypeInt 32 1
+      %SB_RW = OpTypeStruct %int
+%_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
+      %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %19 = OpConstantNull %int
+%atomicAnd_152966 = OpFunction %void None %5
+          %8 = OpLabel
+        %res = OpVariable %_ptr_Function_int Function %19
+         %15 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+          %9 = OpAtomicAnd %int %15 %uint_1 %uint_0 %int_1
+               OpStore %res %9
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+         %21 = OpLabel
+         %22 = OpFunctionCall %void %atomicAnd_152966
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %5
+         %24 = OpLabel
+         %25 = OpFunctionCall %void %atomicAnd_152966
+               OpReturn
+               OpFunctionEnd

test/tint/builtins/atomics/from_gen/literal/atomicAnd/storage_i32.spvasm.expected.glsl

@@ -0,0 +1,68 @@
+#version 310 es
+precision mediump float;
+
+struct SB_RW_atomic {
+  int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  int arg_0;
+} sb_rw;
+void atomicAnd_152966() {
+  int res = 0;
+  int x_9 = atomicAnd(sb_rw.arg_0, 1);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAnd_152966();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct SB_RW_atomic {
+  int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_atomic_1 {
+  int arg_0;
+} sb_rw;
+void atomicAnd_152966() {
+  int res = 0;
+  int x_9 = atomicAnd(sb_rw.arg_0, 1);
+  res = x_9;
+  return;
+}
+
+void compute_main_1() {
+  atomicAnd_152966();
+  return;
+}
+
+void compute_main() {
+  compute_main_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAnd/storage_i32.spvasm.expected.hlsl

@@ -0,0 +1,36 @@
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicAnd(RWByteAddressBuffer buffer, uint offset, int value) {
+  int original_value = 0;
+  buffer.InterlockedAnd(offset, value, original_value);
+  return original_value;
+}
+
+
+void atomicAnd_152966() {
+  int res = 0;
+  const int x_9 = tint_atomicAnd(sb_rw, 0u, 1);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1() {
+  atomicAnd_152966();
+  return;
+}
+
+void fragment_main() {
+  fragment_main_1();
+  return;
+}
+
+void compute_main_1() {
+  atomicAnd_152966();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  compute_main_1();
+  return;
+}

test/tint/builtins/atomics/from_gen/literal/atomicAnd/storage_i32.spvasm.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB_RW_atomic {
+  /* 0x0000 */ atomic_int arg_0;
+};
+
+struct SB_RW {
+  int arg_0;
+};
+
+void atomicAnd_152966(device SB_RW_atomic* const tint_symbol) {
+  int res = 0;
+  int const x_9 = atomic_fetch_and_explicit(&((*(tint_symbol)).arg_0), 1, memory_order_relaxed);
+  res = x_9;
+  return;
+}
+
+void fragment_main_1(device SB_RW_atomic* const tint_symbol_1) {
+  atomicAnd_152966(tint_symbol_1);
+  return;
+}
+
+fragment void fragment_main(device SB_RW_atomic* tint_symbol_2 [[buffer(0)]]) {
+  fragment_main_1(tint_symbol_2);
+  return;
+}
+
+void compute_main_1(device SB_RW_atomic* const tint_symbol_3) {
+  atomicAnd_152966(tint_symbol_3);
+  return;
+}
+
+kernel void compute_main(device SB_RW_atomic* tint_symbol_4 [[buffer(0)]]) {
+  compute_main_1(tint_symbol_4);
+  return;
+}
+