resolver: Enable support of out of order declarations

Allow module-scope declarations to be made in any order.

Bug: tint:1266
Change-Id: Ib2607b6c33fad7c83e2c36f85b0a965eac922ec5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/79769
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/resolver/call_test.cc

@@ -101,6 +101,18 @@ TEST_F(ResolverCallTest, Valid) {
   EXPECT_EQ(call->Target(), Sem().Get(func));
 }
 
+TEST_F(ResolverCallTest, OutOfOrder) {
+  auto* call_expr = Call("b");
+  Func("a", {}, ty.void_(), {CallStmt(call_expr)});
+  auto* b = Func("b", {}, ty.void_(), {});
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* call = Sem().Get(call_expr);
+  EXPECT_NE(call, nullptr);
+  EXPECT_EQ(call->Target(), Sem().Get(b));
+}
+
 }  // namespace
 }  // namespace resolver
 }  // namespace tint

src/resolver/dependency_graph.cc

@@ -474,7 +474,7 @@ struct DependencyAnalysis {
   /// Performs global dependency analysis on the module, emitting any errors to
   /// #diagnostics.
   /// @returns true if analysis found no errors, otherwise false.
-  bool Run(const ast::Module& module, bool allow_out_of_order_decls) {
+  bool Run(const ast::Module& module) {
     // Collect all the named globals from the AST module
     GatherGlobals(module);
 
@@ -487,11 +487,6 @@ struct DependencyAnalysis {
     // Dump the dependency graph if TINT_DUMP_DEPENDENCY_GRAPH is non-zero
     DumpDependencyGraph();
 
-    if (!allow_out_of_order_decls) {
-      // Prevent out-of-order declarations.
-      ErrorOnOutOfOrderDeclarations();
-    }
-
     graph_.ordered_globals = std::move(sorted_);
 
     return !diagnostics_.contains_errors();
@@ -668,36 +663,6 @@ struct DependencyAnalysis {
     return {};
   }
 
-  // TODO(crbug.com/tint/1266): Errors if there are any uses of globals before
-  // their declaration. Out-of-order declarations was added to the WGSL
-  // specification with https://github.com/gpuweb/gpuweb/pull/2244, but Mozilla
-  // have objections to this change so this feature is currently disabled via
-  // this function.
-  void ErrorOnOutOfOrderDeclarations() {
-    if (diagnostics_.contains_errors()) {
-      // Might have already errored about cyclic dependencies. No need to report
-      // out-of-order errors as well.
-      return;
-    }
-    std::unordered_set<const Global*> seen;
-    for (auto* global : declaration_order_) {
-      for (auto* dep : global->deps) {
-        if (!seen.count(dep)) {
-          auto info = DepInfoFor(global, dep);
-          auto name = NameOf(dep->node);
-          AddError(diagnostics_,
-                   KindOf(dep->node) + " '" + name +
-                       "' used before it has been declared",
-                   info.source);
-          AddNote(diagnostics_,
-                  KindOf(dep->node) + " '" + name + "' declared here",
-                  dep->node->source);
-        }
-      }
-      seen.emplace(global);
-    }
-  }
-
   /// CyclicDependencyFound() emits an error diagnostic for a cyclic dependency.
   /// @param root is the global that starts the cyclic dependency, which must be
   /// found in `stack`.
@@ -789,10 +754,9 @@ DependencyGraph::~DependencyGraph() = default;
 bool DependencyGraph::Build(const ast::Module& module,
                             const SymbolTable& symbols,
                             diag::List& diagnostics,
-                            DependencyGraph& output,
+                            DependencyGraph& output) {
-                            bool allow_out_of_order_decls) {
   DependencyAnalysis da{symbols, diagnostics, output};
-  return da.Run(module, allow_out_of_order_decls);
+  return da.Run(module);
 }
 
 }  // namespace resolver

src/resolver/dependency_graph.h

@@ -40,14 +40,11 @@ struct DependencyGraph {
   /// @param symbols the symbol table
   /// @param diagnostics the diagnostic list to populate with errors / warnings
   /// @param output the resulting DependencyGraph
-  /// @param allow_out_of_order_decls if true, then out-of-order declarations
-  /// are not considered an error
   /// @returns true on success, false on error
   static bool Build(const ast::Module& module,
                     const SymbolTable& symbols,
                     diag::List& diagnostics,
-                    DependencyGraph& output,
+                    DependencyGraph& output);
-                    bool allow_out_of_order_decls);
 
   /// All globals in dependency-sorted order.
   std::vector<const ast::Node*> ordered_globals;

src/resolver/dependency_graph_test.cc

@@ -29,13 +29,10 @@ using ::testing::ElementsAre;
 template <typename T>
 class ResolverDependencyGraphTestWithParam : public ResolverTestWithParam<T> {
  public:
-  bool allow_out_of_order_decls = true;
-
   DependencyGraph Build(std::string expected_error = "") {
     DependencyGraph graph;
     auto result = DependencyGraph::Build(this->AST(), this->Symbols(),
-                                         this->Diagnostics(), graph,
+                                         this->Diagnostics(), graph);
-                                         allow_out_of_order_decls);
     if (expected_error.empty()) {
       EXPECT_TRUE(result) << this->Diagnostics().str();
     } else {
@@ -659,13 +656,7 @@ void SymbolTestHelper::Build() {
 ////////////////////////////////////////////////////////////////////////////////
 namespace used_before_decl_tests {
 
-class ResolverDependencyGraphUsedBeforeDeclTest
+using ResolverDependencyGraphUsedBeforeDeclTest = ResolverDependencyGraphTest;
-    : public ResolverDependencyGraphTest {
- public:
-  ResolverDependencyGraphUsedBeforeDeclTest() {
-    allow_out_of_order_decls = false;
-  }
-};
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, FuncCall) {
   // fn A() { B(); }
@@ -674,8 +665,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, FuncCall) {
   Func("A", {}, ty.void_(), {CallStmt(Call(Expr(Source{{12, 34}}, "B")))});
   Func(Source{{56, 78}}, "B", {}, ty.void_(), {Return()});
 
-  Build(R"(12:34 error: function 'B' used before it has been declared
+  Build();
-56:78 note: function 'B' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeConstructed) {
@@ -688,8 +678,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeConstructed) {
        {Block(Ignore(Construct(ty.type_name(Source{{12, 34}}, "T"))))});
   Alias(Source{{56, 78}}, "T", ty.i32());
 
-  Build(R"(12:34 error: alias 'T' used before it has been declared
+  Build();
-56:78 note: alias 'T' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedByLocal) {
@@ -702,8 +691,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedByLocal) {
        {Block(Decl(Var("v", ty.type_name(Source{{12, 34}}, "T"))))});
   Alias(Source{{56, 78}}, "T", ty.i32());
 
-  Build(R"(12:34 error: alias 'T' used before it has been declared
+  Build();
-56:78 note: alias 'T' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedByParam) {
@@ -713,8 +701,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedByParam) {
   Func("F", {Param("p", ty.type_name(Source{{12, 34}}, "T"))}, ty.void_(), {});
   Alias(Source{{56, 78}}, "T", ty.i32());
 
-  Build(R"(12:34 error: alias 'T' used before it has been declared
+  Build();
-56:78 note: alias 'T' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedAsReturnType) {
@@ -724,8 +711,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeUsedAsReturnType) {
   Func("F", {}, ty.type_name(Source{{12, 34}}, "T"), {});
   Alias(Source{{56, 78}}, "T", ty.i32());
 
-  Build(R"(12:34 error: alias 'T' used before it has been declared
+  Build();
-56:78 note: alias 'T' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeByStructMember) {
@@ -735,8 +721,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, TypeByStructMember) {
   Structure("S", {Member("m", ty.type_name(Source{{12, 34}}, "T"))});
   Alias(Source{{56, 78}}, "T", ty.i32());
 
-  Build(R"(12:34 error: alias 'T' used before it has been declared
+  Build();
-56:78 note: alias 'T' declared here)");
 }
 
 TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, VarUsed) {
@@ -751,10 +736,7 @@ TEST_F(ResolverDependencyGraphUsedBeforeDeclTest, VarUsed) {
   Global(Source{{56, 78}}, "G", ty.f32(), ast::StorageClass::kPrivate,
          Expr(2.1f));
 
-  EXPECT_FALSE(r()->Resolve());
+  Build();
-  EXPECT_EQ(r()->error(),
-            R"(12:34 error: var 'G' used before it has been declared
-56:78 note: var 'G' declared here)");
 }
 
 }  // namespace used_before_decl_tests

src/resolver/resolver.cc

@@ -95,8 +95,7 @@ bool Resolver::Resolve() {
   }
 
   if (!DependencyGraph::Build(builder_->AST(), builder_->Symbols(),
-                              builder_->Diagnostics(), dependencies_,
+                              builder_->Diagnostics(), dependencies_)) {
-                              /* allow_out_of_order_decls*/ false)) {
     return false;
   }
 
@@ -118,9 +117,8 @@ bool Resolver::Resolve() {
 bool Resolver::ResolveInternal() {
   Mark(&builder_->AST());
 
-  // Process everything else in the order they appear in the module. This is
+  // Process all module-scope declarations in dependency order.
-  // necessary for validation of use-before-declaration.
+  for (auto* decl : dependencies_.ordered_globals) {
-  for (auto* decl : builder_->AST().GlobalDeclarations()) {
     if (auto* td = decl->As<ast::TypeDecl>()) {
       Mark(td);
       if (!TypeDecl(td)) {

src/transform/add_spirv_block_attribute_test.cc

@@ -548,6 +548,68 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(AddSpirvBlockAttributeTest,
+       Aliases_Nested_OuterBuffer_InnerBuffer_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main() {
+  let f0 = u0.i.f;
+  let f1 = u1.f;
+}
+
+@group(0) @binding(1)
+var<uniform> u1 : MyInner;
+
+type MyInner = Inner;
+
+@group(0) @binding(0)
+var<uniform> u0 : MyOuter;
+
+type MyOuter = Outer;
+
+struct Outer {
+  i : MyInner;
+};
+
+struct Inner {
+  f : f32;
+};
+)";
+  auto* expect = R"(
+@stage(fragment)
+fn main() {
+  let f0 = u0.i.f;
+  let f1 = u1.inner.f;
+}
+
+@internal(spirv_block)
+struct u1_block {
+  inner : Inner;
+}
+
+@group(0) @binding(1) var<uniform> u1 : u1_block;
+
+type MyInner = Inner;
+
+@group(0) @binding(0) var<uniform> u0 : MyOuter;
+
+type MyOuter = Outer;
+
+@internal(spirv_block)
+struct Outer {
+  i : MyInner;
+}
+
+struct Inner {
+  f : f32;
+}
+)";
+
+  auto got = Run<AddSpirvBlockAttribute>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/array_length_from_uniform.cc

@@ -154,10 +154,8 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
       buffer_size_ubo = ctx.dst->Global(
           ctx.dst->Sym(), ctx.dst->ty.Of(buffer_size_struct),
           ast::StorageClass::kUniform,
-          ast::AttributeList{
+          ast::AttributeList{ctx.dst->GroupAndBinding(
-              ctx.dst->create<ast::GroupAttribute>(cfg->ubo_binding.group),
+              cfg->ubo_binding.group, cfg->ubo_binding.binding)});
-              ctx.dst->create<ast::BindingAttribute>(
-                  cfg->ubo_binding.binding)});
     }
     return buffer_size_ubo;
   };

src/transform/array_length_from_uniform_test.cc

@@ -536,6 +536,54 @@ fn main() {
             got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
 }
 
+TEST_F(ArrayLengthFromUniformTest, OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var len : u32 = arrayLength(&sb.arr);
+}
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+struct SB {
+  x : i32;
+  arr : array<i32>;
+};
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  buffer_size : array<vec4<u32>, 1u>;
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+}
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+struct SB {
+  x : i32;
+  arr : array<i32>;
+}
+)";
+
+  ArrayLengthFromUniform::Config cfg({0, 30u});
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
+
+  DataMap data;
+  data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+  EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+            got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/calculate_array_length.cc

@@ -101,7 +101,7 @@ void CalculateArrayLength::Run(CloneContext& ctx,
       auto* type = CreateASTTypeFor(ctx, buffer_type);
       auto* disable_validation = ctx.dst->Disable(
           ast::DisabledValidation::kIgnoreConstructibleFunctionParameter);
-      auto* func = ctx.dst->create<ast::Function>(
+      ctx.dst->AST().AddFunction(ctx.dst->create<ast::Function>(
           name,
           ast::VariableList{
               // Note: The buffer parameter requires the kStorage StorageClass
@@ -118,20 +118,8 @@ void CalculateArrayLength::Run(CloneContext& ctx,
           ast::AttributeList{
               ctx.dst->ASTNodes().Create<BufferSizeIntrinsic>(ctx.dst->ID()),
           },
-          ast::AttributeList{});
+          ast::AttributeList{}));
-      // Insert the intrinsic function after the structure or array structure
+
-      // element type. TODO(crbug.com/tint/1266): Once we allow out-of-order
-      // declarations, this can be simplified.
-      auto* insert_after = buffer_type;
-      while (auto* arr = insert_after->As<sem::Array>()) {
-        insert_after = arr->ElemType();
-      }
-      if (auto* str = insert_after->As<sem::Struct>()) {
-        ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
-                        func);
-      } else {
-        ctx.InsertFront(ctx.src->AST().GlobalDeclarations(), func);
-      }
       return name;
     });
   };

src/transform/calculate_array_length_test.cc

@@ -111,14 +111,14 @@ fn main() {
 )";
 
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
 struct SB {
   x : i32;
   arr : array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> sb : SB;
 
 @stage(compute) @workgroup_size(1)
@@ -149,13 +149,13 @@ fn main() {
 }
 )";
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<S>, result : ptr<function, u32>)
+
 struct S {
   f : f32;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<S>, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> arr : array<S>;
 
 @stage(compute) @workgroup_size(1)
@@ -186,13 +186,13 @@ fn main() {
 }
 )";
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<array<S, 4u>>, result : ptr<function, u32>)
+
 struct S {
   f : f32;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<array<S, 4u>>, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> arr : array<array<S, 4>>;
 
 @stage(compute) @workgroup_size(1)
@@ -261,14 +261,14 @@ fn main() {
 )";
 
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
 struct SB {
   x : i32;
   arr : array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> sb : SB;
 
 @stage(compute) @workgroup_size(1)
@@ -334,15 +334,15 @@ fn main() {
 )";
 
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
 struct SB {
   x : i32;
   y : f32;
   arr : @stride(64) array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> sb : SB;
 
 @stage(compute) @workgroup_size(1)
@@ -381,14 +381,14 @@ fn main() {
 )";
 
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
 struct SB {
   x : i32;
   arr : array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> sb : SB;
 
 @stage(compute) @workgroup_size(1)
@@ -442,6 +442,12 @@ fn main() {
 )";
 
   auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+
 @internal(intrinsic_buffer_size)
 fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
 
@@ -450,17 +456,11 @@ struct SB1 {
   arr1 : array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
-
 struct SB2 {
   x : i32;
   arr2 : array<vec4<f32>>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> sb1 : SB1;
 
 @group(0) @binding(1) var<storage, read> sb2 : SB2;
@@ -512,14 +512,14 @@ fn main() {
 
   auto* expect =
       R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
 struct SB {
   x : i32;
   arr : array<i32>;
 }
 
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
 @group(0) @binding(0) var<storage, read> a : SB;
 
 @group(0) @binding(1) var<storage, read> b : SB;
@@ -544,6 +544,82 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CalculateArrayLengthTest, OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var len1 : u32 = arrayLength(&(sb1.arr1));
+  var len2 : u32 = arrayLength(&(sb2.arr2));
+  var len3 : u32 = arrayLength(&sb3);
+  var x : u32 = (len1 + len2 + len3);
+}
+
+@group(0) @binding(0) var<storage, read> sb1 : SB1;
+
+struct SB1 {
+  x : i32;
+  arr1 : array<i32>;
+};
+
+@group(0) @binding(1) var<storage, read> sb2 : SB2;
+
+struct SB2 {
+  x : i32;
+  arr2 : array<vec4<f32>>;
+};
+
+@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
+)";
+
+  auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var tint_symbol_1 : u32 = 0u;
+  tint_symbol(sb1, &(tint_symbol_1));
+  let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u);
+  var tint_symbol_4 : u32 = 0u;
+  tint_symbol_3(sb2, &(tint_symbol_4));
+  let tint_symbol_5 : u32 = ((tint_symbol_4 - 16u) / 16u);
+  var tint_symbol_7 : u32 = 0u;
+  tint_symbol_6(sb3, &(tint_symbol_7));
+  let tint_symbol_8 : u32 = (tint_symbol_7 / 4u);
+  var len1 : u32 = tint_symbol_2;
+  var len2 : u32 = tint_symbol_5;
+  var len3 : u32 = tint_symbol_8;
+  var x : u32 = ((len1 + len2) + len3);
+}
+
+@group(0) @binding(0) var<storage, read> sb1 : SB1;
+
+struct SB1 {
+  x : i32;
+  arr1 : array<i32>;
+}
+
+@group(0) @binding(1) var<storage, read> sb2 : SB2;
+
+struct SB2 {
+  x : i32;
+  arr2 : array<vec4<f32>>;
+}
+
+@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/combine_samplers_test.cc

@@ -65,6 +65,34 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, SimplePair_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return textureSample(t, s, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var t_s : texture_2d<f32>;
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn main() -> vec4<f32> {
+  return textureSample(t_s, placeholder_sampler, vec2<f32>(1.0, 2.0));
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, SimplePairInAFunction) {
   auto* src = R"(
 @group(0) @binding(0) var t : texture_2d<f32>;
@@ -101,6 +129,42 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, SimplePairInAFunction_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return sample(t, s, vec2<f32>(1.0, 2.0));
+}
+
+fn sample(t : texture_2d<f32>, s : sampler, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t, s, coords);
+}
+
+@group(0) @binding(1) var s : sampler;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var t_s : texture_2d<f32>;
+
+fn main() -> vec4<f32> {
+  return sample(t_s, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn sample(t_s_1 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t_s_1, placeholder_sampler, coords);
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, SimplePairRename) {
   auto* src = R"(
 @group(0) @binding(1) var t : texture_2d<f32>;
@@ -212,6 +276,45 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, AliasedTypes_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return sample(t, s, vec2<f32>(1.0, 2.0));
+}
+
+fn sample(t : Tex2d, s : sampler, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t, s, coords);
+}
+
+@group(0) @binding(0) var t : Tex2d;
+@group(0) @binding(1) var s : sampler;
+
+type Tex2d = texture_2d<f32>;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var t_s : texture_2d<f32>;
+
+fn main() -> vec4<f32> {
+  return sample(t_s, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn sample(t_s_1 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t_s_1, placeholder_sampler, coords);
+}
+
+type Tex2d = texture_2d<f32>;
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, SimplePairInTwoFunctions) {
   auto* src = R"(
 @group(0) @binding(0) var t : texture_2d<f32>;
@@ -256,6 +359,49 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, SimplePairInTwoFunctions_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return f(t, s, vec2<f32>(1.0, 2.0));
+}
+
+fn f(t : texture_2d<f32>, s : sampler, coords : vec2<f32>) -> vec4<f32> {
+  return g(t, s, coords);
+}
+
+fn g(t : texture_2d<f32>, s : sampler, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t, s, coords);
+}
+
+@group(0) @binding(1) var s : sampler;
+@group(0) @binding(0) var t : texture_2d<f32>;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var t_s : texture_2d<f32>;
+
+fn main() -> vec4<f32> {
+  return f(t_s, vec2<f32>(1.0, 2.0));
+}
+
+fn f(t_s_1 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return g(t_s_1, coords);
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn g(t_s_2 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t_s_2, placeholder_sampler, coords);
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, TwoFunctionsGenerateSamePair) {
   auto* src = R"(
 @group(1) @binding(0) var tex : texture_2d<f32>;
@@ -500,6 +646,44 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, GlobalTextureLocalSampler_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return f(samp1, samp2, vec2<f32>(1.0, 2.0));
+}
+
+fn f(s1 : sampler, s2 : sampler, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(tex, s1, coords) + textureSample(tex, s2, coords);
+}
+
+@group(0) @binding(1) var samp1 : sampler;
+@group(0) @binding(2) var samp2 : sampler;
+@group(0) @binding(0) var tex : texture_2d<f32>;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex_samp1 : texture_2d<f32>;
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex_samp2 : texture_2d<f32>;
+
+fn main() -> vec4<f32> {
+  return f(tex_samp1, tex_samp2, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn f(tex_s1 : texture_2d<f32>, tex_s2 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return (textureSample(tex_s1, placeholder_sampler, coords) + textureSample(tex_s2, placeholder_sampler, coords));
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, LocalTextureGlobalSampler) {
   auto* src = R"(
 @group(0) @binding(0) var tex1 : texture_2d<f32>;
@@ -540,6 +724,44 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, LocalTextureGlobalSampler_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return f(tex1, tex2, vec2<f32>(1.0, 2.0));
+}
+
+fn f(t1 : texture_2d<f32>, t2 : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return textureSample(t1, samp, coords) + textureSample(t2, samp, coords);
+}
+
+@group(0) @binding(2) var samp : sampler;
+@group(0) @binding(0) var tex1 : texture_2d<f32>;
+@group(0) @binding(1) var tex2 : texture_2d<f32>;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex1_samp : texture_2d<f32>;
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex2_samp : texture_2d<f32>;
+
+fn main() -> vec4<f32> {
+  return f(tex1_samp, tex2_samp, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn f(t1_samp : texture_2d<f32>, t2_samp : texture_2d<f32>, coords : vec2<f32>) -> vec4<f32> {
+  return (textureSample(t1_samp, placeholder_sampler, coords) + textureSample(t2_samp, placeholder_sampler, coords));
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, TextureLoadNoSampler) {
   auto* src = R"(
 @group(0) @binding(0) var tex : texture_2d<f32>;
@@ -687,6 +909,41 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, TextureSampleCompareInAFunction_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return vec4<f32>(f(tex, samp, vec2<f32>(1.0, 2.0)));
+}
+
+fn f(t : texture_depth_2d, s : sampler_comparison, coords : vec2<f32>) -> f32 {
+  return textureSampleCompare(t, s, coords, 5.0f);
+}
+
+@group(0) @binding(0) var tex : texture_depth_2d;
+@group(0) @binding(1) var samp : sampler_comparison;
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex_samp : texture_depth_2d;
+
+fn main() -> vec4<f32> {
+  return vec4<f32>(f(tex_samp, vec2<f32>(1.0, 2.0)));
+}
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_comparison_sampler : sampler_comparison;
+
+fn f(t_s : texture_depth_2d, coords : vec2<f32>) -> f32 {
+  return textureSampleCompare(t_s, placeholder_comparison_sampler, coords, 5.0);
+}
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CombineSamplersTest, BindingPointCollision) {
   auto* src = R"(
 @group(1) @binding(0) var tex : texture_2d<f32>;
@@ -719,6 +976,37 @@ fn main() -> vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CombineSamplersTest, BindingPointCollision_OutOfOrder) {
+  auto* src = R"(
+fn main() -> vec4<f32> {
+  return textureSample(tex, samp, gcoords);
+}
+
+@group(1) @binding(1) var samp : sampler;
+@group(0) @binding(0) var<uniform> gcoords : vec2<f32>;
+@group(1) @binding(0) var tex : texture_2d<f32>;
+
+)";
+  auto* expect = R"(
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var tex_samp : texture_2d<f32>;
+
+@group(0) @binding(0) @internal(disable_validation__binding_point_collision) var placeholder_sampler : sampler;
+
+fn main() -> vec4<f32> {
+  return textureSample(tex_samp, placeholder_sampler, gcoords);
+}
+
+@internal(disable_validation__binding_point_collision) @group(0) @binding(0) var<uniform> gcoords : vec2<f32>;
+)";
+
+  DataMap data;
+  data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
+                                         sem::BindingPoint());
+  auto got = Run<CombineSamplers>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/external_texture_transform_test.cc

@@ -50,6 +50,34 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ExternalTextureTransformTest, SampleLevelSinglePlane_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleLevel(t, s, (coord.xy / vec2<f32>(4.0, 4.0)));
+}
+
+@group(0) @binding(1) var t : texture_external;
+
+@group(0) @binding(0) var s : sampler;
+)";
+
+  auto* expect = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleLevel(t, s, (coord.xy / vec2<f32>(4.0, 4.0)), 0.0);
+}
+
+@group(0) @binding(1) var t : texture_2d<f32>;
+
+@group(0) @binding(0) var s : sampler;
+)";
+
+  auto got = Run<ExternalTextureTransform>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ExternalTextureTransformTest, LoadSinglePlane) {
   auto* src = R"(
 @group(0) @binding(0) var t : texture_external;
@@ -74,6 +102,30 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ExternalTextureTransformTest, LoadSinglePlane_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureLoad(t, vec2<i32>(1, 1));
+}
+
+@group(0) @binding(0) var t : texture_external;
+)";
+
+  auto* expect = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureLoad(t, vec2<i32>(1, 1), 0);
+}
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+)";
+
+  auto got = Run<ExternalTextureTransform>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ExternalTextureTransformTest, DimensionsSinglePlane) {
   auto* src = R"(
 @group(0) @binding(0) var t : texture_external;
@@ -102,6 +154,34 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ExternalTextureTransformTest, DimensionsSinglePlane_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  var dim : vec2<i32>;
+  dim = textureDimensions(t);
+  return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+}
+
+@group(0) @binding(0) var t : texture_external;
+)";
+
+  auto* expect = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  var dim : vec2<i32>;
+  dim = textureDimensions(t);
+  return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+}
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+)";
+
+  auto got = Run<ExternalTextureTransform>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/first_index_offset_test.cc

@@ -139,6 +139,52 @@ fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32>
   EXPECT_EQ(data->first_instance_offset, 0u);
 }
 
+TEST_F(FirstIndexOffsetTest, BasicModuleVertexIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  test(vert_idx);
+  return vec4<f32>();
+}
+
+fn test(vert_idx : u32) -> u32 {
+  return vert_idx;
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  first_vertex_index : u32;
+}
+
+@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(vertex)
+fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  test((vert_idx + tint_symbol_1.first_vertex_index));
+  return vec4<f32>();
+}
+
+fn test(vert_idx : u32) -> u32 {
+  return vert_idx;
+}
+)";
+
+  DataMap config;
+  config.Add<FirstIndexOffset::BindingPoint>(1, 2);
+  auto got = Run<FirstIndexOffset>(src, std::move(config));
+
+  EXPECT_EQ(expect, str(got));
+
+  auto* data = got.data.Get<FirstIndexOffset::Data>();
+
+  ASSERT_NE(data, nullptr);
+  EXPECT_EQ(data->has_vertex_index, true);
+  EXPECT_EQ(data->has_instance_index, false);
+  EXPECT_EQ(data->first_vertex_offset, 0u);
+  EXPECT_EQ(data->first_instance_offset, 0u);
+}
+
 TEST_F(FirstIndexOffsetTest, BasicModuleInstanceIndex) {
   auto* src = R"(
 fn test(inst_idx : u32) -> u32 {
@@ -185,6 +231,52 @@ fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32
   EXPECT_EQ(data->first_instance_offset, 0u);
 }
 
+TEST_F(FirstIndexOffsetTest, BasicModuleInstanceIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  test(inst_idx);
+  return vec4<f32>();
+}
+
+fn test(inst_idx : u32) -> u32 {
+  return inst_idx;
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  first_instance_index : u32;
+}
+
+@binding(1) @group(7) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(vertex)
+fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  test((inst_idx + tint_symbol_1.first_instance_index));
+  return vec4<f32>();
+}
+
+fn test(inst_idx : u32) -> u32 {
+  return inst_idx;
+}
+)";
+
+  DataMap config;
+  config.Add<FirstIndexOffset::BindingPoint>(1, 7);
+  auto got = Run<FirstIndexOffset>(src, std::move(config));
+
+  EXPECT_EQ(expect, str(got));
+
+  auto* data = got.data.Get<FirstIndexOffset::Data>();
+
+  ASSERT_NE(data, nullptr);
+  EXPECT_EQ(data->has_vertex_index, false);
+  EXPECT_EQ(data->has_instance_index, true);
+  EXPECT_EQ(data->first_vertex_offset, 0u);
+  EXPECT_EQ(data->first_instance_offset, 0u);
+}
+
 TEST_F(FirstIndexOffsetTest, BasicModuleBothIndex) {
   auto* src = R"(
 fn test(instance_idx : u32, vert_idx : u32) -> u32 {
@@ -244,6 +336,65 @@ fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
   EXPECT_EQ(data->first_instance_offset, 4u);
 }
 
+TEST_F(FirstIndexOffsetTest, BasicModuleBothIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
+  test(inputs.instance_idx, inputs.vert_idx);
+  return vec4<f32>();
+}
+
+struct Inputs {
+  @builtin(instance_index) instance_idx : u32;
+  @builtin(vertex_index) vert_idx : u32;
+};
+
+fn test(instance_idx : u32, vert_idx : u32) -> u32 {
+  return instance_idx + vert_idx;
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  first_vertex_index : u32;
+  first_instance_index : u32;
+}
+
+@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(vertex)
+fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
+  test((inputs.instance_idx + tint_symbol_1.first_instance_index), (inputs.vert_idx + tint_symbol_1.first_vertex_index));
+  return vec4<f32>();
+}
+
+struct Inputs {
+  @builtin(instance_index)
+  instance_idx : u32;
+  @builtin(vertex_index)
+  vert_idx : u32;
+}
+
+fn test(instance_idx : u32, vert_idx : u32) -> u32 {
+  return (instance_idx + vert_idx);
+}
+)";
+
+  DataMap config;
+  config.Add<FirstIndexOffset::BindingPoint>(1, 2);
+  auto got = Run<FirstIndexOffset>(src, std::move(config));
+
+  EXPECT_EQ(expect, str(got));
+
+  auto* data = got.data.Get<FirstIndexOffset::Data>();
+
+  ASSERT_NE(data, nullptr);
+  EXPECT_EQ(data->has_vertex_index, true);
+  EXPECT_EQ(data->has_instance_index, true);
+  EXPECT_EQ(data->first_vertex_offset, 0u);
+  EXPECT_EQ(data->first_instance_offset, 4u);
+}
+
 TEST_F(FirstIndexOffsetTest, NestedCalls) {
   auto* src = R"(
 fn func1(vert_idx : u32) -> u32 {
@@ -298,6 +449,60 @@ fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32>
   EXPECT_EQ(data->first_instance_offset, 0u);
 }
 
+TEST_F(FirstIndexOffsetTest, NestedCalls_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  func2(vert_idx);
+  return vec4<f32>();
+}
+
+fn func2(vert_idx : u32) -> u32 {
+  return func1(vert_idx);
+}
+
+fn func1(vert_idx : u32) -> u32 {
+  return vert_idx;
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  first_vertex_index : u32;
+}
+
+@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(vertex)
+fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  func2((vert_idx + tint_symbol_1.first_vertex_index));
+  return vec4<f32>();
+}
+
+fn func2(vert_idx : u32) -> u32 {
+  return func1(vert_idx);
+}
+
+fn func1(vert_idx : u32) -> u32 {
+  return vert_idx;
+}
+)";
+
+  DataMap config;
+  config.Add<FirstIndexOffset::BindingPoint>(1, 2);
+  auto got = Run<FirstIndexOffset>(src, std::move(config));
+
+  EXPECT_EQ(expect, str(got));
+
+  auto* data = got.data.Get<FirstIndexOffset::Data>();
+
+  ASSERT_NE(data, nullptr);
+  EXPECT_EQ(data->has_vertex_index, true);
+  EXPECT_EQ(data->has_instance_index, false);
+  EXPECT_EQ(data->first_vertex_offset, 0u);
+  EXPECT_EQ(data->first_instance_offset, 0u);
+}
+
 TEST_F(FirstIndexOffsetTest, MultipleEntryPoints) {
   auto* src = R"(
 fn func(i : u32) -> u32 {
@@ -369,6 +574,77 @@ fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f
   EXPECT_EQ(data->first_instance_offset, 4u);
 }
 
+TEST_F(FirstIndexOffsetTest, MultipleEntryPoints_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  func(vert_idx);
+  return vec4<f32>();
+}
+
+@stage(vertex)
+fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  func(vert_idx + inst_idx);
+  return vec4<f32>();
+}
+
+@stage(vertex)
+fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  func(inst_idx);
+  return vec4<f32>();
+}
+
+fn func(i : u32) -> u32 {
+  return i;
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  first_vertex_index : u32;
+  first_instance_index : u32;
+}
+
+@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
+
+@stage(vertex)
+fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
+  func((vert_idx + tint_symbol_1.first_vertex_index));
+  return vec4<f32>();
+}
+
+@stage(vertex)
+fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  func(((vert_idx + tint_symbol_1.first_vertex_index) + (inst_idx + tint_symbol_1.first_instance_index)));
+  return vec4<f32>();
+}
+
+@stage(vertex)
+fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
+  func((inst_idx + tint_symbol_1.first_instance_index));
+  return vec4<f32>();
+}
+
+fn func(i : u32) -> u32 {
+  return i;
+}
+)";
+
+  DataMap config;
+  config.Add<FirstIndexOffset::BindingPoint>(1, 2);
+  auto got = Run<FirstIndexOffset>(src, std::move(config));
+
+  EXPECT_EQ(expect, str(got));
+
+  auto* data = got.data.Get<FirstIndexOffset::Data>();
+
+  ASSERT_NE(data, nullptr);
+  EXPECT_EQ(data->has_vertex_index, true);
+  EXPECT_EQ(data->has_instance_index, true);
+  EXPECT_EQ(data->first_vertex_offset, 0u);
+  EXPECT_EQ(data->first_instance_offset, 4u);
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/fold_trivial_single_use_lets_test.cc

@@ -112,6 +112,26 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(FoldTrivialSingleUseLetsTest, NoFold_NonTrivialLet_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  let x = 1;
+  let y = function_with_posssible_side_effect();
+  _ = (x + y);
+}
+
+fn function_with_posssible_side_effect() -> i32 {
+  return 1;
+}
+)";
+
+  auto* expect = src;
+
+  auto got = Run<FoldTrivialSingleUseLets>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(FoldTrivialSingleUseLetsTest, NoFold_UseInSubBlock) {
   auto* src = R"(
 fn f() {

src/transform/localize_struct_array_assignment_test.cc

@@ -90,6 +90,64 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(LocalizeStructArrayAssignmentTest, StructArray_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.a1[uniforms.i] = v;
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+@block struct Uniforms {
+  i : u32;
+};
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  {
+    let tint_symbol = &(s1.a1);
+    var tint_symbol_1 = *(tint_symbol);
+    tint_symbol_1[uniforms.i] = v;
+    *(tint_symbol) = tint_symbol_1;
+  }
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+}
+
+struct InnerS {
+  v : i32;
+}
+
+@block
+struct Uniforms {
+  i : u32;
+}
+)";
+
+  auto got =
+      Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray) {
   auto* src = R"(
 @block struct Uniforms {
@@ -156,6 +214,72 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.s2.a[uniforms.i] = v;
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+struct OuterS {
+  s2 : S1;
+};
+
+struct S1 {
+  a : array<InnerS, 8>;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+@block struct Uniforms {
+  i : u32;
+};
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  {
+    let tint_symbol = &(s1.s2.a);
+    var tint_symbol_1 = *(tint_symbol);
+    tint_symbol_1[uniforms.i] = v;
+    *(tint_symbol) = tint_symbol_1;
+  }
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+struct OuterS {
+  s2 : S1;
+}
+
+struct S1 {
+  a : array<InnerS, 8>;
+}
+
+struct InnerS {
+  v : i32;
+}
+
+@block
+struct Uniforms {
+  i : u32;
+}
+)";
+
+  auto got =
+      Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(LocalizeStructArrayAssignmentTest, StructArrayArray) {
   auto* src = R"(
 @block struct Uniforms {
@@ -437,6 +561,91 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(LocalizeStructArrayAssignmentTest,
+       IndexingWithSideEffectFunc_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s : OuterS;
+  s.a1[getNextIndex()].a2[uniforms.j] = v;
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@block struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+var<private> nextIndex : u32;
+fn getNextIndex() -> u32 {
+  nextIndex = nextIndex + 1u;
+  return nextIndex;
+}
+
+struct OuterS {
+  a1 : array<S1, 8>;
+};
+
+struct S1 {
+  a2 : array<InnerS, 8>;
+};
+
+struct InnerS {
+  v : i32;
+};
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v : InnerS;
+  var s : OuterS;
+  {
+    let tint_symbol = &(s.a1);
+    var tint_symbol_1 = *(tint_symbol);
+    let tint_symbol_2 = &(tint_symbol_1[getNextIndex()].a2);
+    var tint_symbol_3 = *(tint_symbol_2);
+    tint_symbol_3[uniforms.j] = v;
+    *(tint_symbol_2) = tint_symbol_3;
+    *(tint_symbol) = tint_symbol_1;
+  }
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@block
+struct Uniforms {
+  i : u32;
+  j : u32;
+}
+
+var<private> nextIndex : u32;
+
+fn getNextIndex() -> u32 {
+  nextIndex = (nextIndex + 1u);
+  return nextIndex;
+}
+
+struct OuterS {
+  a1 : array<S1, 8>;
+}
+
+struct S1 {
+  a2 : array<InnerS, 8>;
+}
+
+struct InnerS {
+  v : i32;
+}
+)";
+
+  auto got =
+      Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg) {
   auto* src = R"(
 @block struct Uniforms {
@@ -500,6 +709,71 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var s1 : OuterS;
+  f(&s1);
+}
+
+fn f(p : ptr<function, OuterS>) {
+  var v : InnerS;
+  (*p).a1[uniforms.i] = v;
+}
+
+struct InnerS {
+  v : i32;
+};
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@block struct Uniforms {
+  i : u32;
+};
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var s1 : OuterS;
+  f(&(s1));
+}
+
+fn f(p : ptr<function, OuterS>) {
+  var v : InnerS;
+  {
+    let tint_symbol = &((*(p)).a1);
+    var tint_symbol_1 = *(tint_symbol);
+    tint_symbol_1[uniforms.i] = v;
+    *(tint_symbol) = tint_symbol_1;
+  }
+}
+
+struct InnerS {
+  v : i32;
+}
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@block
+struct Uniforms {
+  i : u32;
+}
+)";
+
+  auto got =
+      Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerVar) {
   auto* src = R"(
 @block

src/transform/module_scope_var_to_entry_point_param_test.cc

@@ -63,6 +63,31 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Basic_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  w = p;
+}
+
+var<workgroup> w : f32;
+var<private> p : f32;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  @internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol : f32;
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32;
+  tint_symbol = tint_symbol_1;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, FunctionCalls) {
   auto* src = R"(
 var<private> p : f32;
@@ -116,6 +141,59 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, FunctionCalls_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  foo(1.0);
+}
+
+fn foo(a : f32) {
+  let b : f32 = 2.0;
+  bar(a, b);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32) {
+  p = a;
+  w = b;
+}
+
+var<private> p : f32;
+var<workgroup> w : f32;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
+  @internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_1 : f32;
+  foo(1.0, &(tint_symbol), &(tint_symbol_1));
+}
+
+fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<private, f32>, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<workgroup, f32>) {
+  let b : f32 = 2.0;
+  bar(a, b, tint_symbol_2, tint_symbol_3);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr<private, f32>, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_5 : ptr<workgroup, f32>) {
+  *(tint_symbol_4) = a;
+  *(tint_symbol_5) = b;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Constructors) {
   auto* src = R"(
 var<private> a : f32 = 1.0;
@@ -141,6 +219,31 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Constructors_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let x : f32 = a + b;
+}
+
+var<private> b : f32 = f32();
+var<private> a : f32 = 1.0;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32 = 1.0;
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32 = f32();
+  let x : f32 = (tint_symbol + tint_symbol_1);
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Pointers) {
   auto* src = R"(
 var<private> p : f32;
@@ -172,6 +275,37 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Pointers_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let p_ptr : ptr<private, f32> = &p;
+  let w_ptr : ptr<workgroup, f32> = &w;
+  let x : f32 = *p_ptr + *w_ptr;
+  *p_ptr = x;
+}
+
+var<workgroup> w : f32;
+var<private> p : f32;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
+  @internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_1 : f32;
+  let p_ptr : ptr<private, f32> = &(tint_symbol);
+  let w_ptr : ptr<workgroup, f32> = &(tint_symbol_1);
+  let x : f32 = (*(p_ptr) + *(w_ptr));
+  *(p_ptr) = x;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, FoldAddressOfDeref) {
   auto* src = R"(
 var<private> v : f32;
@@ -211,6 +345,45 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, FoldAddressOfDeref_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  foo();
+}
+
+fn foo() {
+  bar(&v);
+}
+
+fn bar(p : ptr<private, f32>) {
+  (*p) = 0.0;
+}
+
+var<private> v : f32;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  @internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
+  foo(&(tint_symbol));
+}
+
+fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<private, f32>) {
+  bar(tint_symbol_1);
+}
+
+fn bar(p : ptr<private, f32>) {
+  *(p) = 0.0;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_Basic) {
   auto* src = R"(
 struct S {
@@ -246,6 +419,40 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_Basic_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  _ = u;
+  _ = s;
+}
+
+@group(0) @binding(0) var<uniform> u : S;
+@group(0) @binding(1) var<storage> s : S;
+
+struct S {
+  a : f32;
+};
+
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, S>) {
+  _ = *(tint_symbol);
+  _ = *(tint_symbol_1);
+}
+
+struct S {
+  a : f32;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray) {
   auto* src = R"(
 @group(0) @binding(0)
@@ -273,6 +480,33 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  _ = buffer[0];
+}
+
+@group(0) @binding(0)
+var<storage> buffer : array<f32>;
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  arr : array<f32>;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArrayInsideFunction) {
   auto* src = R"(
 @group(0) @binding(0)
@@ -308,6 +542,41 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest,
+       Buffer_RuntimeArrayInsideFunction_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  foo();
+}
+
+fn foo() {
+  _ = buffer[0];
+}
+
+@group(0) @binding(0) var<storage> buffer : array<f32>;
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  arr : array<f32>;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
+  foo(&((*(tint_symbol)).arr));
+}
+
+fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<storage, array<f32>>) {
+  _ = (*(tint_symbol_2))[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_Alias) {
   auto* src = R"(
 type myarray = array<f32>;
@@ -339,6 +608,37 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest,
+       Buffer_RuntimeArray_Alias_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  _ = buffer[0];
+}
+
+@group(0) @binding(0) var<storage> buffer : myarray;
+
+type myarray = array<f32>;
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  arr : array<f32>;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+
+type myarray = array<f32>;
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct) {
   auto* src = R"(
 struct S {
@@ -374,6 +674,40 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  _ = buffer[0];
+}
+
+@group(0) @binding(0) var<storage> buffer : array<S>;
+
+struct S {
+  f : f32;
+};
+)";
+
+  auto* expect = R"(
+struct S {
+  f : f32;
+}
+
+struct tint_symbol_1 {
+  arr : array<S>;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_FunctionCalls) {
   auto* src = R"(
 struct S {
@@ -437,6 +771,69 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_FunctionCalls_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  foo(1.0);
+}
+
+fn foo(a : f32) {
+  let b : f32 = 2.0;
+  _ = u;
+  bar(a, b);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32) {
+  _ = u;
+  _ = s;
+}
+
+struct S {
+  a : f32;
+};
+
+@group(0) @binding(0)
+var<uniform> u : S;
+@group(0) @binding(1)
+var<storage> s : S;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, S>) {
+  foo(1.0, tint_symbol, tint_symbol_1);
+}
+
+fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S>) {
+  let b : f32 = 2.0;
+  _ = *(tint_symbol_2);
+  bar(a, b, tint_symbol_2, tint_symbol_3);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr<uniform, S>, @internal(disable_validation__ignore_storage_class) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_5 : ptr<storage, S>) {
+  _ = *(tint_symbol_4);
+  _ = *(tint_symbol_5);
+}
+
+struct S {
+  a : f32;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, HandleTypes_Basic) {
   auto* src = R"(
 @group(0) @binding(0) var t : texture_2d<f32>;
@@ -515,6 +912,60 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest,
+       HandleTypes_FunctionCalls_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  foo(1.0);
+}
+
+fn foo(a : f32) {
+  let b : f32 = 2.0;
+  _ = t;
+  bar(a, b);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32) {
+  _ = t;
+  _ = s;
+}
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0) @binding(1) var s : sampler;
+)";
+
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol : texture_2d<f32>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_1 : sampler) {
+  foo(1.0, tint_symbol, tint_symbol_1);
+}
+
+fn foo(a : f32, tint_symbol_2 : texture_2d<f32>, tint_symbol_3 : sampler) {
+  let b : f32 = 2.0;
+  _ = tint_symbol_2;
+  bar(a, b, tint_symbol_2, tint_symbol_3);
+  no_uses();
+}
+
+fn no_uses() {
+}
+
+fn bar(a : f32, b : f32, tint_symbol_4 : texture_2d<f32>, tint_symbol_5 : sampler) {
+  _ = tint_symbol_4;
+  _ = tint_symbol_5;
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Matrix) {
   auto* src = R"(
 var<workgroup> m : mat2x2<f32>;
@@ -626,6 +1077,49 @@ fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr
   EXPECT_EQ(expect, str(got));
 }
 
+// Test that we do not duplicate a struct type used by multiple workgroup
+// variables that are promoted to threadgroup memory arguments.
+TEST_F(ModuleScopeVarToEntryPointParamTest,
+       DuplicateThreadgroupArgumentTypes_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+  let x = a;
+  let y = b;
+}
+
+var<workgroup> a : S;
+var<workgroup> b : S;
+
+struct S {
+  m : mat2x2<f32>;
+};
+)";
+
+  auto* expect = R"(
+struct S {
+  m : mat2x2<f32>;
+}
+
+struct tint_symbol_3 {
+  a : S;
+  b : S;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
+  let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
+  let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
+  let x = *(tint_symbol);
+  let y = *(tint_symbol_2);
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, UnusedVariables) {
   auto* src = R"(
 struct S {

src/transform/multiplanar_external_texture_test.cc

@@ -136,6 +136,49 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the transform works with a textureDimensions call.
+TEST_F(MultiplanarExternalTextureTest, Dimensions_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  var dim : vec2<i32>;
+  dim = textureDimensions(ext_tex);
+  return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+}
+
+@group(0) @binding(0) var ext_tex : texture_external;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(1) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(2) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  var dim : vec2<i32>;
+  dim = textureDimensions(ext_tex);
+  return vec4<f32>(0.0, 0.0, 0.0, 0.0);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+)";
+
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{{{0, 0}, {{0, 1}, {0, 2}}}});
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Test that the transform works with a textureSampleLevel call.
 TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleLevel) {
   auto* src = R"(
@@ -192,6 +235,62 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+// Test that the transform works with a textureSampleLevel call.
+TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleLevel_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleLevel(ext_tex, s, coord.xy);
+}
+
+@group(0) @binding(1) var ext_tex : texture_external;
+@group(0) @binding(0) var s : sampler;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params);
+}
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(0) var s : sampler;
+)";
+
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with a textureLoad call.
 TEST_F(MultiplanarExternalTextureTest, BasicTextureLoad) {
   auto* src = R"(
@@ -245,6 +344,59 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the transform works with a textureLoad call.
+TEST_F(MultiplanarExternalTextureTest, BasicTextureLoad_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureLoad(ext_tex, vec2<i32>(1, 1));
+}
+
+@group(0) @binding(0) var ext_tex : texture_external;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(1) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(2) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureLoad(plane0, coord, 0);
+  }
+  let y = (textureLoad(plane0, coord, 0).r - 0.0625);
+  let uv = (textureLoad(plane1, coord, 0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+)";
+
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{{{0, 0}, {{0, 1}, {0, 2}}}});
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with both a textureSampleLevel and textureLoad
 // call.
 TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad) {
@@ -316,6 +468,77 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the transform works with both a textureSampleLevel and textureLoad
+// call.
+TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return textureSampleLevel(ext_tex, s, coord.xy) + textureLoad(ext_tex, vec2<i32>(1, 1));
+}
+
+@group(0) @binding(0) var s : sampler;
+@group(0) @binding(1) var ext_tex : texture_external;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureLoad(plane0, coord, 0);
+  }
+  let y = (textureLoad(plane0, coord, 0).r - 0.0625);
+  let uv = (textureLoad(plane1, coord, 0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+@stage(fragment)
+fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
+  return (textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params));
+}
+
+@group(0) @binding(0) var s : sampler;
+
+@group(0) @binding(1) var ext_tex : texture_2d<f32>;
+)";
+
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{{{0, 1}, {{0, 2}, {0, 3}}}});
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with many instances of texture_external.
 TEST_F(MultiplanarExternalTextureTest, ManyTextureSampleLevel) {
   auto* src = R"(
@@ -464,6 +687,73 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the texture_external passed as a function parameter produces the
+// correct output.
+TEST_F(MultiplanarExternalTextureTest,
+       ExternalTexturePassedAsParam_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main() {
+  f(ext_tex, smp);
+}
+
+fn f(t : texture_external, s : sampler) {
+  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) var ext_tex : texture_external;
+@group(0) @binding(1) var smp : sampler;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@stage(fragment)
+fn main() {
+  f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
+  textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(1) var smp : sampler;
+)";
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{
+          {{0, 0}, {{0, 2}, {0, 3}}},
+      });
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the texture_external passed as a parameter not in the first
 // position produces the correct output.
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsSecondParam) {
@@ -606,6 +896,84 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that multiple texture_external params passed to a function produces the
+// correct output.
+TEST_F(MultiplanarExternalTextureTest,
+       ExternalTexturePassedAsParamMultiple_OutOfOrder) {
+  auto* src = R"(
+@stage(fragment)
+fn main() {
+  f(ext_tex, smp, ext_tex2);
+}
+
+fn f(t : texture_external, s : sampler, t2 : texture_external) {
+  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+  textureSampleLevel(t2, s, vec2<f32>(1.0, 2.0));
+}
+
+@group(0) @binding(0) var ext_tex : texture_external;
+@group(0) @binding(1) var smp : sampler;
+@group(0) @binding(2) var ext_tex2 : texture_external;
+
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(3) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(4) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@group(0) @binding(5) var ext_tex_plane_1_1 : texture_2d<f32>;
+
+@group(0) @binding(6) var<uniform> ext_tex_params_1 : ExternalTextureParams;
+
+@stage(fragment)
+fn main() {
+  f(ext_tex, ext_tex_plane_1, ext_tex_params, smp, ext_tex2, ext_tex_plane_1_1, ext_tex_params_1);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler, t2 : texture_2d<f32>, ext_tex_plane_1_3 : texture_2d<f32>, ext_tex_params_3 : ExternalTextureParams) {
+  textureSampleExternal(t, ext_tex_plane_1_2, s, vec2<f32>(1.0, 2.0), ext_tex_params_2);
+  textureSampleExternal(t2, ext_tex_plane_1_3, s, vec2<f32>(1.0, 2.0), ext_tex_params_3);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(1) var smp : sampler;
+
+@group(0) @binding(2) var ext_tex2 : texture_2d<f32>;
+)";
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{
+          {{0, 0}, {{0, 3}, {0, 4}}},
+          {{0, 2}, {{0, 5}, {0, 6}}},
+      });
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the texture_external passed to as a parameter to multiple
 // functions produces the correct output.
 TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamNested) {
@@ -680,6 +1048,81 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the texture_external passed to as a parameter to multiple
+// functions produces the correct output.
+TEST_F(MultiplanarExternalTextureTest,
+       ExternalTexturePassedAsParamNested_OutOfOrder) {
+  auto* src = R"(
+fn nested(t : texture_external, s : sampler) {
+  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+}
+
+fn f(t : texture_external, s : sampler) {
+  nested(t, s);
+}
+
+@group(0) @binding(0) var ext_tex : texture_external;
+@group(0) @binding(1) var smp : sampler;
+
+@stage(fragment)
+fn main() {
+  f(ext_tex, smp);
+}
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+fn nested(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
+  textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
+}
+
+fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2 : ExternalTextureParams, s : sampler) {
+  nested(t, ext_tex_plane_1_2, ext_tex_params_2, s);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(1) var smp : sampler;
+
+@stage(fragment)
+fn main() {
+  f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
+}
+)";
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{
+          {{0, 0}, {{0, 2}, {0, 3}}},
+      });
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 // Tests that the transform works with a function using an external texture,
 // even if there's no external texture declared at module scope.
 TEST_F(MultiplanarExternalTextureTest,
@@ -780,6 +1223,75 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Tests that the the transform handles aliases to external textures
+TEST_F(MultiplanarExternalTextureTest, ExternalTextureAlias_OutOfOrder) {
+  auto* src = R"(
+[[stage(fragment)]]
+fn main() {
+  f(ext_tex, smp);
+}
+
+fn f(t : ET, s : sampler) {
+  textureSampleLevel(t, s, vec2<f32>(1.0, 2.0));
+}
+
+[[group(0), binding(0)]] var ext_tex : ET;
+[[group(0), binding(1)]] var smp : sampler;
+
+type ET = texture_external;
+)";
+
+  auto* expect = R"(
+struct ExternalTextureParams {
+  numPlanes : u32;
+  vr : f32;
+  ug : f32;
+  vg : f32;
+  ub : f32;
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@stage(fragment)
+fn main() {
+  f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
+}
+
+fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp : sampler, coord : vec2<f32>, params : ExternalTextureParams) -> vec4<f32> {
+  if ((params.numPlanes == 1u)) {
+    return textureSampleLevel(plane0, smp, coord, 0.0);
+  }
+  let y = (textureSampleLevel(plane0, smp, coord, 0.0).r - 0.0625);
+  let uv = (textureSampleLevel(plane1, smp, coord, 0.0).rg - 0.5);
+  let u = uv.x;
+  let v = uv.y;
+  let r = ((1.164000034 * y) + (params.vr * v));
+  let g = (((1.164000034 * y) - (params.ug * u)) - (params.vg * v));
+  let b = ((1.164000034 * y) + (params.ub * u));
+  return vec4<f32>(r, g, b, 1.0);
+}
+
+fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1 : ExternalTextureParams, s : sampler) {
+  textureSampleExternal(t, ext_tex_plane_1_1, s, vec2<f32>(1.0, 2.0), ext_tex_params_1);
+}
+
+@group(0) @binding(0) var ext_tex : texture_2d<f32>;
+
+@group(0) @binding(1) var smp : sampler;
+
+type ET = texture_external;
+)";
+  DataMap data;
+  data.Add<MultiplanarExternalTexture::NewBindingPoints>(
+      MultiplanarExternalTexture::BindingsMap{
+          {{0, 0}, {{0, 2}, {0, 3}}},
+      });
+  auto got = Run<MultiplanarExternalTexture>(src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/num_workgroups_from_uniform_test.cc

@@ -146,6 +146,52 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(NumWorkgroupsFromUniformTest, StructOnlyMember_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main(in : Builtins) {
+  let groups_x = in.num_wgs.x;
+  let groups_y = in.num_wgs.y;
+  let groups_z = in.num_wgs.z;
+}
+
+struct Builtins {
+  @builtin(num_workgroups) num_wgs : vec3<u32>;
+};
+)";
+
+  auto* expect = R"(
+struct tint_symbol_2 {
+  num_workgroups : vec3<u32>;
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+fn main_inner(in : Builtins) {
+  let groups_x = in.num_wgs.x;
+  let groups_y = in.num_wgs.y;
+  let groups_z = in.num_wgs.z;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  main_inner(Builtins(tint_symbol_3.num_workgroups));
+}
+
+struct Builtins {
+  num_wgs : vec3<u32>;
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+  data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(NumWorkgroupsFromUniformTest, StructMultipleMembers) {
   auto* src = R"(
 struct Builtins {
@@ -203,6 +249,64 @@ fn main(tint_symbol : tint_symbol_1) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(NumWorkgroupsFromUniformTest, StructMultipleMembers_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main(in : Builtins) {
+  let groups_x = in.num_wgs.x;
+  let groups_y = in.num_wgs.y;
+  let groups_z = in.num_wgs.z;
+}
+
+struct Builtins {
+  @builtin(global_invocation_id) gid : vec3<u32>;
+  @builtin(num_workgroups) num_wgs : vec3<u32>;
+  @builtin(workgroup_id) wgid : vec3<u32>;
+};
+
+)";
+
+  auto* expect = R"(
+struct tint_symbol_2 {
+  num_workgroups : vec3<u32>;
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+struct tint_symbol_1 {
+  @builtin(global_invocation_id)
+  gid : vec3<u32>;
+  @builtin(workgroup_id)
+  wgid : vec3<u32>;
+}
+
+fn main_inner(in : Builtins) {
+  let groups_x = in.num_wgs.x;
+  let groups_y = in.num_wgs.y;
+  let groups_z = in.num_wgs.z;
+}
+
+@stage(compute) @workgroup_size(1)
+fn main(tint_symbol : tint_symbol_1) {
+  main_inner(Builtins(tint_symbol.gid, tint_symbol_3.num_workgroups, tint_symbol.wgid));
+}
+
+struct Builtins {
+  gid : vec3<u32>;
+  num_wgs : vec3<u32>;
+  wgid : vec3<u32>;
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+  data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(NumWorkgroupsFromUniformTest, MultipleEntryPoints) {
   auto* src = R"(
 struct Builtins1 {

src/transform/pad_array_elements_test.cc

@@ -216,6 +216,42 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PadArrayElementsTest, ArrayAlias_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var arr : Array;
+  arr = Array();
+  arr = Array(1, 2, 3, 4);
+  let vals : Array = Array(1, 2, 3, 4);
+  arr = vals;
+  let x = arr[3];
+}
+
+type Array = @stride(16) array<i32, 4>;
+)";
+  auto* expect = R"(
+struct tint_padded_array_element {
+  @size(16)
+  el : i32;
+}
+
+fn f() {
+  var arr : array<tint_padded_array_element, 4u>;
+  arr = array<tint_padded_array_element, 4u>();
+  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
+  let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
+  arr = vals;
+  let x = arr[3].el;
+}
+
+type Array = array<tint_padded_array_element, 4u>;
+)";
+
+  auto got = Run<PadArrayElements>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PadArrayElementsTest, ArraysInStruct) {
   auto* src = R"(
 struct S {
@@ -364,6 +400,65 @@ fn f(s : S) -> i32 {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct_OutOfOrder) {
+  auto* src = R"(
+fn f(s : S) -> i32 {
+  return s.a[2] + s.b[1][2] + s.c[3][1][2];
+}
+
+struct S {
+  a : @stride(512) array<i32, 4>;
+  b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
+  c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
+};
+)";
+  auto* expect = R"(
+struct tint_padded_array_element {
+  @size(512)
+  el : i32;
+}
+
+struct tint_padded_array_element_1 {
+  @size(32)
+  el : i32;
+}
+
+struct tint_padded_array_element_2 {
+  @size(512)
+  el : array<tint_padded_array_element_1, 4u>;
+}
+
+struct tint_padded_array_element_3 {
+  @size(8)
+  el : i32;
+}
+
+struct tint_padded_array_element_4 {
+  @size(64)
+  el : array<tint_padded_array_element_3, 4u>;
+}
+
+struct tint_padded_array_element_5 {
+  @size(512)
+  el : array<tint_padded_array_element_4, 4u>;
+}
+
+fn f(s : S) -> i32 {
+  return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
+}
+
+struct S {
+  a : array<tint_padded_array_element, 4u>;
+  b : array<tint_padded_array_element_2, 4u>;
+  c : array<tint_padded_array_element_5, 4u>;
+}
+)";
+
+  auto got = Run<PadArrayElements>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PadArrayElementsTest, DeclarationOrder) {
   auto* src = R"(
 type T0 = i32;

src/transform/promote_initializers_to_const_var_test.cc

@@ -91,6 +91,38 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PromoteInitializersToConstVarTest, BasicStruct_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var x = S(1, 2.0, vec3<f32>()).b;
+}
+
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+};
+)";
+
+  auto* expect = R"(
+fn f() {
+  let tint_symbol = S(1, 2.0, vec3<f32>());
+  var x = tint_symbol.b;
+}
+
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+}
+)";
+
+  DataMap data;
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopInit) {
   auto* src = R"(
 fn f() {
@@ -155,6 +187,44 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PromoteInitializersToConstVarTest, StructInForLoopInit_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var insert_after = 1;
+  for(var x = S(1, 2.0, vec3<f32>()).b; ; ) {
+    break;
+  }
+}
+
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+};
+)";
+
+  auto* expect = R"(
+fn f() {
+  var insert_after = 1;
+  let tint_symbol = S(1, 2.0, vec3<f32>());
+  for(var x = tint_symbol.b; ; ) {
+    break;
+  }
+}
+
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+}
+)";
+
+  DataMap data;
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopCond) {
   auto* src = R"(
 fn f() {
@@ -460,6 +530,46 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PromoteInitializersToConstVarTest, Mixed_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var x = S2(array<S1, 3u>(S1(1), S1(2), S1(3))).a[1].a;
+}
+
+struct S2 {
+  a : array<S1, 3u>;
+};
+
+struct S1 {
+  a : i32;
+};
+)";
+
+  auto* expect = R"(
+fn f() {
+  let tint_symbol = S1(1);
+  let tint_symbol_1 = S1(2);
+  let tint_symbol_2 = S1(3);
+  let tint_symbol_3 = array<S1, 3u>(tint_symbol, tint_symbol_1, tint_symbol_2);
+  let tint_symbol_4 = S2(tint_symbol_3);
+  var x = tint_symbol_4.a[1].a;
+}
+
+struct S2 {
+  a : array<S1, 3u>;
+}
+
+struct S1 {
+  a : i32;
+}
+)";
+
+  DataMap data;
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PromoteInitializersToConstVarTest, NoChangeOnVarDecl) {
   auto* src = R"(
 struct S {
@@ -486,6 +596,32 @@ let module_str : S = S(1, 2.0, 3);
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PromoteInitializersToConstVarTest, NoChangeOnVarDecl_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var local_arr = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
+  var local_str = S(1, 2.0, 3);
+}
+
+let module_str : S = S(1, 2.0, 3);
+
+struct S {
+  a : i32;
+  b : f32;
+  c : i32;
+}
+
+let module_arr : array<f32, 4u> = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
+)";
+
+  auto* expect = src;
+
+  DataMap data;
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/remove_phonies_test.cc

@@ -130,6 +130,54 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RemovePhoniesTest, SingleSideEffects_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  _ = neg(1);
+  _ = add(2, 3);
+  _ = add(neg(4), neg(5));
+  _ = u32(neg(6));
+  _ = f32(add(7, 8));
+  _ = vec2<f32>(f32(neg(9)));
+  _ = vec3<i32>(1, neg(10), 3);
+  _ = mat2x2<f32>(1.0, f32(add(11, 12)), 3.0, 4.0);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+  return (a + b);
+}
+
+fn neg(a : i32) -> i32 {
+  return -(a);
+}
+)";
+
+  auto* expect = R"(
+fn f() {
+  neg(1);
+  add(2, 3);
+  add(neg(4), neg(5));
+  neg(6);
+  add(7, 8);
+  neg(9);
+  neg(10);
+  add(11, 12);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+  return (a + b);
+}
+
+fn neg(a : i32) -> i32 {
+  return -(a);
+}
+)";
+
+  auto got = Run<RemovePhonies>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RemovePhoniesTest, MultipleSideEffects) {
   auto* src = R"(
 fn neg(a : i32) -> i32 {
@@ -187,6 +235,63 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RemovePhoniesTest, MultipleSideEffects_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  _ = (1 + add(2 + add(3, 4), 5)) * add(6, 7) * neg(8);
+  _ = add(9, neg(10)) + neg(11);
+  _ = xor(12u, 13u) + xor(14u, 15u);
+  _ = neg(16) / neg(17) + add(18, 19);
+}
+
+fn neg(a : i32) -> i32 {
+  return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+  return (a + b);
+}
+
+fn xor(a : u32, b : u32) -> u32 {
+  return (a ^ b);
+}
+)";
+
+  auto* expect = R"(
+fn phony_sink(p0 : i32, p1 : i32, p2 : i32) {
+}
+
+fn phony_sink_1(p0 : i32, p1 : i32) {
+}
+
+fn phony_sink_2(p0 : u32, p1 : u32) {
+}
+
+fn f() {
+  phony_sink(add((2 + add(3, 4)), 5), add(6, 7), neg(8));
+  phony_sink_1(add(9, neg(10)), neg(11));
+  phony_sink_2(xor(12u, 13u), xor(14u, 15u));
+  phony_sink(neg(16), neg(17), add(18, 19));
+}
+
+fn neg(a : i32) -> i32 {
+  return -(a);
+}
+
+fn add(a : i32, b : i32) -> i32 {
+  return (a + b);
+}
+
+fn xor(a : u32, b : u32) -> u32 {
+  return (a ^ b);
+}
+)";
+
+  auto got = Run<RemovePhonies>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RemovePhoniesTest, ForLoop) {
   auto* src = R"(
 struct S {
@@ -254,6 +359,73 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RemovePhoniesTest, ForLoop_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  for (_ = &s.arr; ;_ = &s.arr) {
+    break;
+  }
+  for (_ = x(); ;_ = y() + z()) {
+    break;
+  }
+}
+
+fn x() -> i32 {
+  return 0;
+}
+
+fn y() -> i32 {
+  return 0;
+}
+
+fn z() -> i32 {
+  return 0;
+}
+
+struct S {
+  arr : array<i32>;
+};
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+)";
+
+  auto* expect = R"(
+fn phony_sink(p0 : i32, p1 : i32) {
+}
+
+fn f() {
+  for(; ; ) {
+    break;
+  }
+  for(x(); ; phony_sink(y(), z())) {
+    break;
+  }
+}
+
+fn x() -> i32 {
+  return 0;
+}
+
+fn y() -> i32 {
+  return 0;
+}
+
+fn z() -> i32 {
+  return 0;
+}
+
+struct S {
+  arr : array<i32>;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+)";
+
+  auto got = Run<RemovePhonies>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/robustness_test.cc

@@ -48,6 +48,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_Clamp_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  let b : f32 = a[c];
+}
+
+let c : u32 = 1u;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  let b : f32 = a[1u];
+}
+
+let c : u32 = 1u;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Array_Idx_Nested_Scalar) {
   auto* src = R"(
 var<private> a : array<f32, 3>;
@@ -78,6 +104,36 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_Nested_Scalar_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var c : f32 = a[ b[i] ];
+}
+
+var<private> i : u32;
+
+var<private> b : array<i32, 5>;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)];
+}
+
+var<private> i : u32;
+
+var<private> b : array<i32, 5>;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Array_Idx_Scalar) {
   auto* src = R"(
 var<private> a : array<f32, 3>;
@@ -100,6 +156,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_Scalar_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[1];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[1];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Array_Idx_Expr) {
   auto* src = R"(
 var<private> a : array<f32, 3>;
@@ -126,6 +204,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_Expr_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[c + 2 - 3];
+}
+
+var<private> c : i32;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
+}
+
+var<private> c : i32;
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Array_Idx_Negative) {
   auto* src = R"(
 var<private> a : array<f32, 3>;
@@ -148,6 +252,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_Negative_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[-1];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[0];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Array_Idx_OutOfBounds) {
   auto* src = R"(
 var<private> a : array<f32, 3>;
@@ -170,6 +296,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Array_Idx_OutOfBounds_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[3];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2];
+}
+
+var<private> a : array<f32, 3>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 // TODO(crbug.com/tint/1177) - Validation currently forbids arrays larger than
 // 0xffffffff. If WGSL supports 64-bit indexing, re-enable this test.
 TEST_F(RobustnessTest, DISABLED_LargeArrays_Idx) {
@@ -272,6 +420,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Idx_Scalar_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[1];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[1];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Idx_Expr) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -298,6 +468,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Idx_Expr_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[c + 2 - 3];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[min(u32(((c + 2) - 3)), 2u)];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Swizzle_Idx_Scalar) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -320,6 +516,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Swizzle_Idx_Scalar_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a.xy[2];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a.xy[1];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Swizzle_Idx_Var) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -345,6 +563,33 @@ fn f() {
 
   EXPECT_EQ(expect, str(got));
 }
+
+TEST_F(RobustnessTest, Vector_Swizzle_Idx_Var_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a.xy[c];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a.xy[min(u32(c), 1u)];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Swizzle_Idx_Expr) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -371,6 +616,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Swizzle_Idx_Expr_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a.xy[c + 2 - 3];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)];
+}
+
+var<private> c : i32;
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Idx_Negative) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -393,6 +664,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Idx_Negative_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[-1];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[0];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Vector_Idx_OutOfBounds) {
   auto* src = R"(
 var<private> a : vec3<f32>;
@@ -415,6 +708,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Vector_Idx_OutOfBounds_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[3];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2];
+}
+
+var<private> a : vec3<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_Scalar) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -437,6 +752,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_Scalar_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[2][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_Expr_Column) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -463,6 +800,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_Expr_Column_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[c + 2 - 3][1];
+}
+
+var<private> c : i32;
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1];
+}
+
+var<private> c : i32;
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_Expr_Row) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -489,6 +852,32 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_Expr_Row_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[1][c + 2 - 3];
+}
+
+var<private> c : i32;
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)];
+}
+
+var<private> c : i32;
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_Negative_Column) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -511,6 +900,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_Negative_Column_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[-1][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[0][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_Negative_Row) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -533,6 +944,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_Negative_Row_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[2][-1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2][0];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_OutOfBounds_Column) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -555,6 +988,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_OutOfBounds_Column_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[5][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(RobustnessTest, Matrix_Idx_OutOfBounds_Row) {
   auto* src = R"(
 var<private> a : mat3x2<f32>;
@@ -577,6 +1032,28 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(RobustnessTest, Matrix_Idx_OutOfBounds_Row_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var b : f32 = a[2][5];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto* expect = R"(
+fn f() {
+  var b : f32 = a[2][1];
+}
+
+var<private> a : mat3x2<f32>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 // TODO(dsinclair): Implement when constant_id exists
 TEST_F(RobustnessTest, DISABLED_Vector_Constant_Id_Clamps) {
   // @override(1300) let idx : i32;
@@ -644,9 +1121,36 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
-// TODO(dsinclair): Clamp atomics when available.
+TEST_F(RobustnessTest, RuntimeArray_Clamps_OutOfOrder) {
-TEST_F(RobustnessTest, DISABLED_Atomics_Clamp) {
+  auto* src = R"(
-  FAIL();
+fn f() {
+  var d : f32 = s.b[25];
+}
+
+@group(0) @binding(0) var<storage, read> s : S;
+
+struct S {
+  a : f32;
+  b : array<f32>;
+};
+)";
+
+  auto* expect = R"(
+fn f() {
+  var d : f32 = s.b[min(25u, (arrayLength(&(s.b)) - 1u))];
+}
+
+@group(0) @binding(0) var<storage, read> s : S;
+
+struct S {
+  a : f32;
+  b : array<f32>;
+}
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
 }
 
 // Clamp textureLoad() coord, array_index and level values
@@ -715,6 +1219,72 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Clamp textureLoad() coord, array_index and level values
+TEST_F(RobustnessTest, TextureLoad_Clamp_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var array_idx : i32;
+  var level_idx : i32;
+  var sample_idx : i32;
+
+  textureLoad(tex_1d, 1, level_idx);
+  textureLoad(tex_2d, vec2<i32>(1, 2), level_idx);
+  textureLoad(tex_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
+  textureLoad(tex_3d, vec3<i32>(1, 2, 3), level_idx);
+  textureLoad(tex_ms_2d, vec2<i32>(1, 2), sample_idx);
+  textureLoad(tex_depth_2d, vec2<i32>(1, 2), level_idx);
+  textureLoad(tex_depth_2d_arr, vec2<i32>(1, 2), array_idx, level_idx);
+  textureLoad(tex_external, vec2<i32>(1, 2));
+}
+
+@group(0) @binding(0) var tex_1d : texture_1d<f32>;
+@group(0) @binding(0) var tex_2d : texture_2d<f32>;
+@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
+@group(0) @binding(0) var tex_3d : texture_3d<f32>;
+@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
+@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
+@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
+@group(0) @binding(0) var tex_external : texture_external;
+)";
+
+  auto* expect =
+      R"(
+fn f() {
+  var array_idx : i32;
+  var level_idx : i32;
+  var sample_idx : i32;
+  textureLoad(tex_1d, clamp(1, i32(), (textureDimensions(tex_1d, clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1))) - i32(1))), clamp(level_idx, 0, (textureNumLevels(tex_1d) - 1)));
+  textureLoad(tex_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d, clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_2d) - 1)));
+  textureLoad(tex_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_2d_arr) - 1)));
+  textureLoad(tex_3d, clamp(vec3<i32>(1, 2, 3), vec3<i32>(), (textureDimensions(tex_3d, clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1))) - vec3<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_3d) - 1)));
+  textureLoad(tex_ms_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_ms_2d) - vec2<i32>(1))), sample_idx);
+  textureLoad(tex_depth_2d, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1))) - vec2<i32>(1))), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d) - 1)));
+  textureLoad(tex_depth_2d_arr, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_depth_2d_arr, clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1))) - vec2<i32>(1))), clamp(array_idx, 0, (textureNumLayers(tex_depth_2d_arr) - 1)), clamp(level_idx, 0, (textureNumLevels(tex_depth_2d_arr) - 1)));
+  textureLoad(tex_external, clamp(vec2<i32>(1, 2), vec2<i32>(), (textureDimensions(tex_external) - vec2<i32>(1))));
+}
+
+@group(0) @binding(0) var tex_1d : texture_1d<f32>;
+
+@group(0) @binding(0) var tex_2d : texture_2d<f32>;
+
+@group(0) @binding(0) var tex_2d_arr : texture_2d_array<f32>;
+
+@group(0) @binding(0) var tex_3d : texture_3d<f32>;
+
+@group(0) @binding(0) var tex_ms_2d : texture_multisampled_2d<f32>;
+
+@group(0) @binding(0) var tex_depth_2d : texture_depth_2d;
+
+@group(0) @binding(0) var tex_depth_2d_arr : texture_depth_2d_array;
+
+@group(0) @binding(0) var tex_external : texture_external;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 // Clamp textureStore() coord, array_index and level values
 TEST_F(RobustnessTest, TextureStore_Clamp) {
   auto* src = R"(
@@ -756,6 +1326,48 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+// Clamp textureStore() coord, array_index and level values
+TEST_F(RobustnessTest, TextureStore_Clamp_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  textureStore(tex1d, 10, vec4<i32>());
+  textureStore(tex2d, vec2<i32>(10, 20), vec4<i32>());
+  textureStore(tex2d_arr, vec2<i32>(10, 20), 50, vec4<i32>());
+  textureStore(tex3d, vec3<i32>(10, 20, 30), vec4<i32>());
+}
+
+@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
+
+@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
+
+@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
+
+@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
+
+)";
+
+  auto* expect = R"(
+fn f() {
+  textureStore(tex1d, clamp(10, i32(), (textureDimensions(tex1d) - i32(1))), vec4<i32>());
+  textureStore(tex2d, clamp(vec2<i32>(10, 20), vec2<i32>(), (textureDimensions(tex2d) - vec2<i32>(1))), vec4<i32>());
+  textureStore(tex2d_arr, clamp(vec2<i32>(10, 20), vec2<i32>(), (textureDimensions(tex2d_arr) - vec2<i32>(1))), clamp(50, 0, (textureNumLayers(tex2d_arr) - 1)), vec4<i32>());
+  textureStore(tex3d, clamp(vec3<i32>(10, 20, 30), vec3<i32>(), (textureDimensions(tex3d) - vec3<i32>(1))), vec4<i32>());
+}
+
+@group(0) @binding(0) var tex1d : texture_storage_1d<rgba8sint, write>;
+
+@group(0) @binding(1) var tex2d : texture_storage_2d<rgba8sint, write>;
+
+@group(0) @binding(2) var tex2d_arr : texture_storage_2d_array<rgba8sint, write>;
+
+@group(0) @binding(3) var tex3d : texture_storage_3d<rgba8sint, write>;
+)";
+
+  auto got = Run<Robustness>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 // TODO(dsinclair): Test for scoped variables when shadowing is implemented
 TEST_F(RobustnessTest, DISABLED_Shadowed_Variable) {
   // var a : array<f32, 3>;

src/transform/unshadow_test.cc

@@ -84,6 +84,36 @@ fn Y() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, LocalShadowsAlias_OutOfOrder) {
+  auto* src = R"(
+fn X() {
+  var a = false;
+}
+
+fn Y() {
+  let a = true;
+}
+
+type a = i32;
+)";
+
+  auto* expect = R"(
+fn X() {
+  var a_1 = false;
+}
+
+fn Y() {
+  let a_2 = true;
+}
+
+type a = i32;
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(UnshadowTest, LocalShadowsStruct) {
   auto* src = R"(
 struct a {
@@ -118,13 +148,50 @@ fn Y() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, LocalShadowsStruct_OutOfOrder) {
+  auto* src = R"(
+fn X() {
+  var a = true;
+}
+
+fn Y() {
+  let a = false;
+}
+
+struct a {
+  m : i32;
+};
+
+)";
+
+  auto* expect = R"(
+fn X() {
+  var a_1 = true;
+}
+
+fn Y() {
+  let a_2 = false;
+}
+
+struct a {
+  m : i32;
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(UnshadowTest, LocalShadowsFunction) {
   auto* src = R"(
 fn a() {
   var a = true;
+  var b = false;
 }
 
 fn b() {
+  let a = true;
   let b = false;
 }
 )";
@@ -132,11 +199,44 @@ fn b() {
   auto* expect = R"(
 fn a() {
   var a_1 = true;
+  var b_1 = false;
 }
 
 fn b() {
+  let a_2 = true;
+  let b_2 = false;
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsFunction_OutOfOrder) {
+  auto* src = R"(
+fn b() {
+  let a = true;
+  let b = false;
+}
+
+fn a() {
+  var a = true;
+  var b = false;
+}
+
+)";
+
+  auto* expect = R"(
+fn b() {
+  let a_1 = true;
   let b_1 = false;
 }
+
+fn a() {
+  var a_2 = true;
+  var b_2 = false;
+}
 )";
 
   auto got = Run<Unshadow>(src);
@@ -174,6 +274,36 @@ fn Y() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, LocalShadowsGlobalVar_OutOfOrder) {
+  auto* src = R"(
+fn X() {
+  var a = (a == 123);
+}
+
+fn Y() {
+  let a = (a == 321);
+}
+
+var<private> a : i32;
+)";
+
+  auto* expect = R"(
+fn X() {
+  var a_1 = (a == 123);
+}
+
+fn Y() {
+  let a_2 = (a == 321);
+}
+
+var<private> a : i32;
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(UnshadowTest, LocalShadowsGlobalLet) {
   auto* src = R"(
 let a : i32 = 1;
@@ -204,6 +334,36 @@ fn Y() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, LocalShadowsGlobalLet_OutOfOrder) {
+  auto* src = R"(
+fn X() {
+  var a = (a == 123);
+}
+
+fn Y() {
+  let a = (a == 321);
+}
+
+let a : i32 = 1;
+)";
+
+  auto* expect = R"(
+fn X() {
+  var a_1 = (a == 123);
+}
+
+fn Y() {
+  let a_2 = (a == 321);
+}
+
+let a : i32 = 1;
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(UnshadowTest, LocalShadowsLocalVar) {
   auto* src = R"(
 fn X() {
@@ -360,6 +520,26 @@ fn F(a_1 : bool) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, ParamShadowsGlobalLet_OutOfOrder) {
+  auto* src = R"(
+fn F(a : bool) {
+}
+
+let a : i32 = 1;
+)";
+
+  auto* expect = R"(
+fn F(a_1 : bool) {
+}
+
+let a : i32 = 1;
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(UnshadowTest, ParamShadowsAlias) {
   auto* src = R"(
 type a = i32;
@@ -392,6 +572,38 @@ fn F(a_1 : a) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(UnshadowTest, ParamShadowsAlias_OutOfOrder) {
+  auto* src = R"(
+fn F(a : a) {
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+
+type a = i32;
+)";
+
+  auto* expect = R"(
+fn F(a_1 : a) {
+  {
+    var a_2 = (a_1 == 123);
+  }
+  {
+    let a_3 = (a_1 == 321);
+  }
+}
+
+type a = i32;
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/vertex_pulling_test.cc

@@ -426,6 +426,86 @@ fn main(tint_symbol_1 : tint_symbol) -> @builtin(position) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(VertexPullingTest,
+       ExistingVertexIndexAndInstanceIndex_Struct_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
+  return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
+}
+
+struct Inputs {
+  @location(0) var_a : f32;
+  @location(1) var_b : f32;
+  @builtin(vertex_index) custom_vertex_index : u32;
+  @builtin(instance_index) custom_instance_index : u32;
+};
+)";
+
+  auto* expect = R"(
+struct TintVertexData {
+  tint_vertex_data : @stride(4) array<u32>;
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@binding(1) @group(4) var<storage, read> tint_pulling_vertex_buffer_1 : TintVertexData;
+
+struct tint_symbol {
+  @builtin(vertex_index)
+  custom_vertex_index : u32;
+  @builtin(instance_index)
+  custom_instance_index : u32;
+}
+
+@stage(vertex)
+fn main(tint_symbol_1 : tint_symbol) -> @builtin(position) vec4<f32> {
+  var inputs : Inputs;
+  inputs.custom_vertex_index = tint_symbol_1.custom_vertex_index;
+  inputs.custom_instance_index = tint_symbol_1.custom_instance_index;
+  {
+    let buffer_array_base_0 = inputs.custom_vertex_index;
+    inputs.var_a = bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[buffer_array_base_0]);
+    let buffer_array_base_1 = inputs.custom_instance_index;
+    inputs.var_b = bitcast<f32>(tint_pulling_vertex_buffer_1.tint_vertex_data[buffer_array_base_1]);
+  }
+  return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
+}
+
+struct Inputs {
+  @location(0)
+  var_a : f32;
+  @location(1)
+  var_b : f32;
+  @builtin(vertex_index)
+  custom_vertex_index : u32;
+  @builtin(instance_index)
+  custom_instance_index : u32;
+}
+)";
+
+  VertexPulling::Config cfg;
+  cfg.vertex_state = {{
+      {
+          4,
+          VertexStepMode::kVertex,
+          {{VertexFormat::kFloat32, 0, 0}},
+      },
+      {
+          4,
+          VertexStepMode::kInstance,
+          {{VertexFormat::kFloat32, 0, 1}},
+      },
+  }};
+  cfg.entry_point_name = "main";
+
+  DataMap data;
+  data.Add<VertexPulling::Config>(cfg);
+  auto got = Run<VertexPulling>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(VertexPullingTest, ExistingVertexIndexAndInstanceIndex_SeparateStruct) {
   auto* src = R"(
 struct Inputs {
@@ -502,6 +582,83 @@ fn main(indices : Indices) -> @builtin(position) vec4<f32> {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(VertexPullingTest,
+       ExistingVertexIndexAndInstanceIndex_SeparateStruct_OutOfOrder) {
+  auto* src = R"(
+@stage(vertex)
+fn main(inputs : Inputs, indices : Indices) -> @builtin(position) vec4<f32> {
+  return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
+}
+
+struct Inputs {
+  @location(0) var_a : f32;
+  @location(1) var_b : f32;
+};
+
+struct Indices {
+  @builtin(vertex_index) custom_vertex_index : u32;
+  @builtin(instance_index) custom_instance_index : u32;
+};
+)";
+
+  auto* expect = R"(
+struct TintVertexData {
+  tint_vertex_data : @stride(4) array<u32>;
+}
+
+@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
+
+@binding(1) @group(4) var<storage, read> tint_pulling_vertex_buffer_1 : TintVertexData;
+
+@stage(vertex)
+fn main(indices : Indices) -> @builtin(position) vec4<f32> {
+  var inputs : Inputs;
+  {
+    let buffer_array_base_0 = indices.custom_vertex_index;
+    inputs.var_a = bitcast<f32>(tint_pulling_vertex_buffer_0.tint_vertex_data[buffer_array_base_0]);
+    let buffer_array_base_1 = indices.custom_instance_index;
+    inputs.var_b = bitcast<f32>(tint_pulling_vertex_buffer_1.tint_vertex_data[buffer_array_base_1]);
+  }
+  return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
+}
+
+struct Inputs {
+  @location(0)
+  var_a : f32;
+  @location(1)
+  var_b : f32;
+}
+
+struct Indices {
+  @builtin(vertex_index)
+  custom_vertex_index : u32;
+  @builtin(instance_index)
+  custom_instance_index : u32;
+}
+)";
+
+  VertexPulling::Config cfg;
+  cfg.vertex_state = {{
+      {
+          4,
+          VertexStepMode::kVertex,
+          {{VertexFormat::kFloat32, 0, 0}},
+      },
+      {
+          4,
+          VertexStepMode::kInstance,
+          {{VertexFormat::kFloat32, 0, 1}},
+      },
+  }};
+  cfg.entry_point_name = "main";
+
+  DataMap data;
+  data.Add<VertexPulling::Config>(cfg);
+  auto got = Run<VertexPulling>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(VertexPullingTest, TwoAttributesSameBuffer) {
   auto* src = R"(
 @stage(vertex)

src/transform/wrap_arrays_in_structs_test.cc

@@ -172,6 +172,48 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(WrapArraysInStructsTest, ArrayAlias_OutOfOrder) {
+  auto* src = R"(
+fn f() {
+  var arr : Array;
+  arr = Array();
+  arr = Array(Inner(1, 2), Inner(3, 4));
+  let vals : Array = Array(Inner(1, 2), Inner(3, 4));
+  arr = vals;
+  let x = arr[3];
+}
+
+type Array = array<Inner, 2>;
+type Inner = array<i32, 2>;
+)";
+  auto* expect = R"(
+struct tint_array_wrapper_1 {
+  arr : array<i32, 2u>;
+}
+
+struct tint_array_wrapper {
+  arr : array<tint_array_wrapper_1, 2u>;
+}
+
+fn f() {
+  var arr : tint_array_wrapper;
+  arr = tint_array_wrapper(array<tint_array_wrapper_1, 2u>());
+  arr = tint_array_wrapper(array<tint_array_wrapper_1, 2u>(tint_array_wrapper_1(array<i32, 2u>(1, 2)), tint_array_wrapper_1(array<i32, 2u>(3, 4))));
+  let vals : tint_array_wrapper = tint_array_wrapper(array<tint_array_wrapper_1, 2u>(tint_array_wrapper_1(array<i32, 2u>(1, 2)), tint_array_wrapper_1(array<i32, 2u>(3, 4))));
+  arr = vals;
+  let x = arr.arr[3];
+}
+
+type Array = tint_array_wrapper;
+
+type Inner = tint_array_wrapper_1;
+)";
+
+  auto got = Run<WrapArraysInStructs>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(WrapArraysInStructsTest, ArraysInStruct) {
   auto* src = R"(
 struct S {
@@ -326,6 +368,57 @@ fn f2() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(WrapArraysInStructsTest, DeclarationOrder_OutOfOrder) {
+  auto* src = R"(
+fn f2() {
+  var v : array<i32, 3>;
+}
+
+type T3 = i32;
+
+fn f1(a : array<i32, 2>) {
+}
+
+type T2 = i32;
+
+type T1 = array<i32, 1>;
+
+type T0 = i32;
+)";
+  auto* expect = R"(
+struct tint_array_wrapper {
+  arr : array<i32, 3u>;
+}
+
+fn f2() {
+  var v : tint_array_wrapper;
+}
+
+type T3 = i32;
+
+struct tint_array_wrapper_1 {
+  arr : array<i32, 2u>;
+}
+
+fn f1(a : tint_array_wrapper_1) {
+}
+
+type T2 = i32;
+
+struct tint_array_wrapper_2 {
+  arr : array<i32, 1u>;
+}
+
+type T1 = tint_array_wrapper_2;
+
+type T0 = i32;
+)";
+
+  auto got = Run<WrapArraysInStructs>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/zero_init_workgroup_memory_test.cc

@@ -93,6 +93,29 @@ fn f() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, UnreferencedWorkgroupVars_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+}
+
+fn unreferenced() {
+  b = c;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : i32;
+
+var<workgroup> c : i32;
+)";
+  auto* expect = src;
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndex) {
   auto* src = R"(
 var<workgroup> v : i32;
@@ -120,6 +143,34 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       SingleWorkgroupVar_ExistingLocalIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  _ = v; // Initialization should be inserted above this statement
+}
+
+var<workgroup> v : i32;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  {
+    v = i32();
+  }
+  workgroupBarrier();
+  _ = v;
+}
+
+var<workgroup> v : i32;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest,
        SingleWorkgroupVar_ExistingLocalIndexInStruct) {
   auto* src = R"(
@@ -157,6 +208,43 @@ fn f(params : Params) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       SingleWorkgroupVar_ExistingLocalIndexInStruct_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f(params : Params) {
+  _ = v; // Initialization should be inserted above this statement
+}
+
+struct Params {
+  @builtin(local_invocation_index) local_idx : u32;
+};
+
+var<workgroup> v : i32;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(params : Params) {
+  {
+    v = i32();
+  }
+  workgroupBarrier();
+  _ = v;
+}
+
+struct Params {
+  @builtin(local_invocation_index)
+  local_idx : u32;
+}
+
+var<workgroup> v : i32;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_InjectedLocalIndex) {
   auto* src = R"(
 var<workgroup> v : i32;
@@ -184,6 +272,34 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       SingleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+  _ = v; // Initialization should be inserted above this statement
+}
+
+var<workgroup> v : i32;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
+  {
+    v = i32();
+  }
+  workgroupBarrier();
+  _ = v;
+}
+
+var<workgroup> v : i32;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest,
        MultipleWorkgroupVar_ExistingLocalIndex_Size1) {
   auto* src = R"(
@@ -248,6 +364,70 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       MultipleWorkgroupVar_ExistingLocalIndex_Size1_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  _ = a; // Initialization should be inserted above this statement
+  _ = b;
+  _ = c;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+};
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  {
+    a = i32();
+    b.x = i32();
+  }
+  for(var idx : u32 = local_idx; (idx < 8u); idx = (idx + 1u)) {
+    let i : u32 = idx;
+    b.y[i] = i32();
+  }
+  for(var idx_1 : u32 = local_idx; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
+    let i_1 : u32 = idx_1;
+    c[i_1].x = i32();
+  }
+  for(var idx_2 : u32 = local_idx; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
+    let i_2 : u32 = (idx_2 / 8u);
+    let i : u32 = (idx_2 % 8u);
+    c[i_2].y[i] = i32();
+  }
+  workgroupBarrier();
+  _ = a;
+  _ = b;
+  _ = c;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+}
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest,
        MultipleWorkgroupVar_ExistingLocalIndex_Size_2_3) {
   auto* src = R"(
@@ -534,6 +714,70 @@ fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(loc
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       MultipleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
+  _ = a; // Initialization should be inserted above this statement
+  _ = b;
+  _ = c;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+};
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index : u32) {
+  {
+    a = i32();
+    b.x = i32();
+  }
+  for(var idx : u32 = local_invocation_index; (idx < 8u); idx = (idx + 1u)) {
+    let i : u32 = idx;
+    b.y[i] = i32();
+  }
+  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
+    let i_1 : u32 = idx_1;
+    c[i_1].x = i32();
+  }
+  for(var idx_2 : u32 = local_invocation_index; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
+    let i_2 : u32 = (idx_2 / 8u);
+    let i : u32 = (idx_2 % 8u);
+    c[i_2].y[i] = i32();
+  }
+  workgroupBarrier();
+  _ = a;
+  _ = b;
+  _ = c;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+}
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_MultipleEntryPoints) {
   auto* src = R"(
 struct S {
@@ -633,6 +877,106 @@ fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       MultipleWorkgroupVar_MultipleEntryPoints_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f1() {
+  _ = a; // Initialization should be inserted above this statement
+  _ = c;
+}
+
+@stage(compute) @workgroup_size(1, 2, 3)
+fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
+  _ = b; // Initialization should be inserted above this statement
+}
+
+@stage(compute) @workgroup_size(4, 5, 6)
+fn f3() {
+  _ = c; // Initialization should be inserted above this statement
+  _ = a;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+};
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
+  {
+    a = i32();
+  }
+  for(var idx : u32 = local_invocation_index; (idx < 32u); idx = (idx + 1u)) {
+    let i : u32 = idx;
+    c[i].x = i32();
+  }
+  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 256u); idx_1 = (idx_1 + 1u)) {
+    let i_1 : u32 = (idx_1 / 8u);
+    let i_2 : u32 = (idx_1 % 8u);
+    c[i_1].y[i_2] = i32();
+  }
+  workgroupBarrier();
+  _ = a;
+  _ = c;
+}
+
+@stage(compute) @workgroup_size(1, 2, 3)
+fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index_1 : u32) {
+  if ((local_invocation_index_1 < 1u)) {
+    b.x = i32();
+  }
+  for(var idx_2 : u32 = local_invocation_index_1; (idx_2 < 8u); idx_2 = (idx_2 + 6u)) {
+    let i_3 : u32 = idx_2;
+    b.y[i_3] = i32();
+  }
+  workgroupBarrier();
+  _ = b;
+}
+
+@stage(compute) @workgroup_size(4, 5, 6)
+fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
+  if ((local_invocation_index_2 < 1u)) {
+    a = i32();
+  }
+  if ((local_invocation_index_2 < 32u)) {
+    let i_4 : u32 = local_invocation_index_2;
+    c[i_4].x = i32();
+  }
+  for(var idx_3 : u32 = local_invocation_index_2; (idx_3 < 256u); idx_3 = (idx_3 + 120u)) {
+    let i_5 : u32 = (idx_3 / 8u);
+    let i_6 : u32 = (idx_3 % 8u);
+    c[i_5].y[i_6] = i32();
+  }
+  workgroupBarrier();
+  _ = c;
+  _ = a;
+}
+
+var<workgroup> a : i32;
+
+var<workgroup> b : S;
+
+var<workgroup> c : array<S, 32>;
+
+struct S {
+  x : i32;
+  y : array<i32, 8>;
+}
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage) {
   auto* src = R"(
 var<workgroup> v : i32;
@@ -676,6 +1020,49 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  call_use_v(); // Initialization should be inserted above this statement
+}
+
+fn call_use_v() {
+  use_v();
+}
+
+fn use_v() {
+  _ = v;
+}
+
+var<workgroup> v : i32;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_idx : u32) {
+  {
+    v = i32();
+  }
+  workgroupBarrier();
+  call_use_v();
+}
+
+fn call_use_v() {
+  use_v();
+}
+
+fn use_v() {
+  _ = v;
+}
+
+var<workgroup> v : i32;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics) {
   auto* src = R"(
 var<workgroup> i : atomic<i32>;
@@ -709,6 +1096,39 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+  atomicLoad(&(i)); // Initialization should be inserted above this statement
+  atomicLoad(&(u));
+}
+
+var<workgroup> i : atomic<i32>;
+var<workgroup> u : atomic<u32>;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
+  {
+    atomicStore(&(i), i32());
+    atomicStore(&(u), u32());
+  }
+  workgroupBarrier();
+  atomicLoad(&(i));
+  atomicLoad(&(u));
+}
+
+var<workgroup> i : atomic<i32>;
+
+var<workgroup> u : atomic<u32>;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupStructOfAtomics) {
   auto* src = R"(
 struct S {
@@ -756,6 +1176,53 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupStructOfAtomics_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+  _ = w.a; // Initialization should be inserted above this statement
+}
+
+var<workgroup> w : S;
+
+struct S {
+  a : i32;
+  i : atomic<i32>;
+  b : f32;
+  u : atomic<u32>;
+  c : u32;
+};
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
+  {
+    w.a = i32();
+    atomicStore(&(w.i), i32());
+    w.b = f32();
+    atomicStore(&(w.u), u32());
+    w.c = u32();
+  }
+  workgroupBarrier();
+  _ = w.a;
+}
+
+var<workgroup> w : S;
+
+struct S {
+  a : i32;
+  i : atomic<i32>;
+  b : f32;
+  u : atomic<u32>;
+  c : u32;
+}
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics) {
   auto* src = R"(
 var<workgroup> w : array<atomic<u32>, 4>;
@@ -784,6 +1251,34 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+  atomicLoad(&w[0]); // Initialization should be inserted above this statement
+}
+
+var<workgroup> w : array<atomic<u32>, 4>;
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
+  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+    let i : u32 = idx;
+    atomicStore(&(w[i]), u32());
+  }
+  workgroupBarrier();
+  atomicLoad(&(w[0]));
+}
+
+var<workgroup> w : array<atomic<u32>, 4>;
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfStructOfAtomics) {
   auto* src = R"(
 struct S {
@@ -832,6 +1327,55 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest,
+       WorkgroupArrayOfStructOfAtomics_OutOfOrder) {
+  auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn f() {
+  _ = w[0].a; // Initialization should be inserted above this statement
+}
+
+var<workgroup> w : array<S, 4>;
+
+struct S {
+  a : i32;
+  i : atomic<i32>;
+  b : f32;
+  u : atomic<u32>;
+  c : u32;
+};
+)";
+  auto* expect = R"(
+@stage(compute) @workgroup_size(1)
+fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
+  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+    let i_1 : u32 = idx;
+    w[i_1].a = i32();
+    atomicStore(&(w[i_1].i), i32());
+    w[i_1].b = f32();
+    atomicStore(&(w[i_1].u), u32());
+    w[i_1].c = u32();
+  }
+  workgroupBarrier();
+  _ = w[0].a;
+}
+
+var<workgroup> w : array<S, 4>;
+
+struct S {
+  a : i32;
+  i : atomic<i32>;
+  b : f32;
+  u : atomic<u32>;
+  c : u32;
+}
+)";
+
+  auto got = Run<ZeroInitWorkgroupMemory>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/writer/hlsl/generator_impl_sanitizer_test.cc

@@ -181,7 +181,6 @@ TEST_F(HlslSanitizerTest, Call_ArrayLength_ArrayLengthFromUniform) {
   auto* expect = R"(cbuffer cbuffer_tint_symbol_1 : register(b4, space3) {
   uint4 tint_symbol_1[2];
 };
-
 ByteAddressBuffer b : register(t1, space2);
 ByteAddressBuffer c : register(t2, space2);
 

test/out_of_order_decls/alias/alias.wgsl

@@ -0,0 +1,7 @@
+type T1 = T2;
+type T2 = i32;
+
+@stage(fragment)
+fn f() {
+  var v : T1;
+}

test/out_of_order_decls/alias/alias.wgsl.expected.glsl

@@ -0,0 +1,11 @@
+#version 310 es
+precision mediump float;
+
+void f() {
+  int v = 0;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/alias/alias.wgsl.expected.hlsl

@@ -0,0 +1,4 @@
+void f() {
+  int v = 0;
+  return;
+}

test/out_of_order_decls/alias/alias.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+fragment void f() {
+  int v = 0;
+  return;
+}
+

test/out_of_order_decls/alias/alias.wgsl.expected.spvasm

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %f "f"
+               OpName %v "v"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+          %8 = OpConstantNull %int
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_int Function %8
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/alias/alias.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+type T1 = T2;
+
+type T2 = i32;
+
+@stage(fragment)
+fn f() {
+  var v : T1;
+}

test/out_of_order_decls/alias/struct.wgsl

@@ -0,0 +1,10 @@
+type T = S;
+
+struct S {
+  m : i32;
+}
+
+@stage(fragment)
+fn f() {
+  var v : T;
+}

test/out_of_order_decls/alias/struct.wgsl.expected.glsl

@@ -0,0 +1,15 @@
+#version 310 es
+precision mediump float;
+
+struct S {
+  int m;
+};
+
+void f() {
+  S v = S(0);
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/alias/struct.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+struct S {
+  int m;
+};
+
+void f() {
+  S v = (S)0;
+  return;
+}

test/out_of_order_decls/alias/struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int m;
+};
+
+fragment void f() {
+  S v = {};
+  return;
+}
+

test/out_of_order_decls/alias/struct.wgsl.expected.spvasm

@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %f "f"
+               OpName %S "S"
+               OpMemberName %S 0 "m"
+               OpName %v "v"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+%_ptr_Function_S = OpTypePointer Function %S
+          %9 = OpConstantNull %S
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_S Function %9
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/alias/struct.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+type T = S;
+
+struct S {
+  m : i32;
+}
+
+@stage(fragment)
+fn f() {
+  var v : T;
+}

test/out_of_order_decls/array/alias.wgsl

@@ -0,0 +1,7 @@
+var<private> A : array<T, 4>;
+type T = i32;
+
+@stage(fragment)
+fn f() {
+  A[0] = 1;
+}

test/out_of_order_decls/array/alias.wgsl.expected.glsl

@@ -0,0 +1,12 @@
+#version 310 es
+precision mediump float;
+
+int A[4] = int[4](0, 0, 0, 0);
+void f() {
+  A[0] = 1;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/array/alias.wgsl.expected.hlsl

@@ -0,0 +1,6 @@
+static int A[4] = (int[4])0;
+
+void f() {
+  A[0] = 1;
+  return;
+}

test/out_of_order_decls/array/alias.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[4];
+};
+
+fragment void f() {
+  thread tint_array_wrapper tint_symbol = {};
+  tint_symbol.arr[0] = 1;
+  return;
+}
+

test/out_of_order_decls/array/alias.wgsl.expected.spvasm

@@ -0,0 +1,30 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %A "A"
+               OpName %f "f"
+               OpDecorate %_arr_int_uint_4 ArrayStride 4
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_int_uint_4 = OpTypeArray %int %uint_4
+%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
+          %7 = OpConstantNull %_arr_int_uint_4
+          %A = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Private_int = OpTypePointer Private %int
+      %int_1 = OpConstant %int 1
+          %f = OpFunction %void None %8
+         %11 = OpLabel
+         %14 = OpAccessChain %_ptr_Private_int %A %int_0
+               OpStore %14 %int_1
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/array/alias.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+var<private> A : array<T, 4>;
+
+type T = i32;
+
+@stage(fragment)
+fn f() {
+  A[0] = 1;
+}

test/out_of_order_decls/array/struct.wgsl

@@ -0,0 +1,9 @@
+var<private> A : array<S, 4>;
+struct S {
+  m : i32;
+};
+
+@stage(fragment)
+fn f() {
+  A[0] = S(1);
+}

test/out_of_order_decls/array/struct.wgsl.expected.glsl

@@ -0,0 +1,17 @@
+#version 310 es
+precision mediump float;
+
+struct S {
+  int m;
+};
+
+S A[4] = S[4](S(0), S(0), S(0), S(0));
+void f() {
+  S tint_symbol = S(1);
+  A[0] = tint_symbol;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/array/struct.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+struct S {
+  int m;
+};
+
+static S A[4] = (S[4])0;
+
+void f() {
+  const S tint_symbol = {1};
+  A[0] = tint_symbol;
+  return;
+}

test/out_of_order_decls/array/struct.wgsl.expected.msl

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int m;
+};
+
+struct tint_array_wrapper {
+  S arr[4];
+};
+
+fragment void f() {
+  thread tint_array_wrapper tint_symbol_1 = {};
+  S const tint_symbol = {.m=1};
+  tint_symbol_1.arr[0] = tint_symbol;
+  return;
+}
+

test/out_of_order_decls/array/struct.wgsl.expected.spvasm

@@ -0,0 +1,35 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %S "S"
+               OpMemberName %S 0 "m"
+               OpName %A "A"
+               OpName %f "f"
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_arr_S_uint_4 ArrayStride 4
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_S_uint_4 = OpTypeArray %S %uint_4
+%_ptr_Private__arr_S_uint_4 = OpTypePointer Private %_arr_S_uint_4
+          %8 = OpConstantNull %_arr_S_uint_4
+          %A = OpVariable %_ptr_Private__arr_S_uint_4 Private %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+      %int_0 = OpConstant %int 0
+%_ptr_Private_S = OpTypePointer Private %S
+      %int_1 = OpConstant %int 1
+         %17 = OpConstantComposite %S %int_1
+          %f = OpFunction %void None %9
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Private_S %A %int_0
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/array/struct.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+var<private> A : array<S, 4>;
+
+struct S {
+  m : i32;
+}
+
+@stage(fragment)
+fn f() {
+  A[0] = S(1);
+}

test/out_of_order_decls/func/func.wgsl

@@ -0,0 +1,7 @@
+@stage(fragment)
+fn f1() {
+  f2();
+}
+
+fn f2() {
+}

test/out_of_order_decls/func/func.wgsl.expected.glsl

@@ -0,0 +1,14 @@
+#version 310 es
+precision mediump float;
+
+void f2() {
+}
+
+void f1() {
+  f2();
+}
+
+void main() {
+  f1();
+  return;
+}

test/out_of_order_decls/func/func.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+void f2() {
+}
+
+void f1() {
+  f2();
+  return;
+}

test/out_of_order_decls/func/func.wgsl.expected.msl

@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void f2() {
+}
+
+fragment void f1() {
+  f2();
+  return;
+}
+

test/out_of_order_decls/func/func.wgsl.expected.spvasm

@@ -0,0 +1,22 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 8
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f1 "f1"
+               OpExecutionMode %f1 OriginUpperLeft
+               OpName %f2 "f2"
+               OpName %f1 "f1"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+         %f2 = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+         %f1 = OpFunction %void None %1
+          %6 = OpLabel
+          %7 = OpFunctionCall %void %f2
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/func/func.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+@stage(fragment)
+fn f1() {
+  f2();
+}
+
+fn f2() {
+}

test/out_of_order_decls/func/let.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  let b = a;
+}
+
+let a : i32 = 1;

test/out_of_order_decls/func/let.wgsl.expected.glsl

@@ -0,0 +1,12 @@
+#version 310 es
+precision mediump float;
+
+const int a = 1;
+void f() {
+  int b = a;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/func/let.wgsl.expected.hlsl

@@ -0,0 +1,6 @@
+static const int a = 1;
+
+void f() {
+  const int b = a;
+  return;
+}

test/out_of_order_decls/func/let.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+constant int a = 1;
+
+fragment void f() {
+  int const b = a;
+  return;
+}
+

test/out_of_order_decls/func/let.wgsl.expected.spvasm

@@ -0,0 +1,19 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 7
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %a "a"
+               OpName %f "f"
+        %int = OpTypeInt 32 1
+          %a = OpConstant %int 1
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+          %f = OpFunction %void None %3
+          %6 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/func/let.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  let b = a;
+}
+
+let a : i32 = 1;

test/out_of_order_decls/func/type.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  var b : T;
+}
+
+type T = i32;

test/out_of_order_decls/func/type.wgsl.expected.glsl

@@ -0,0 +1,11 @@
+#version 310 es
+precision mediump float;
+
+void f() {
+  int b = 0;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/func/type.wgsl.expected.hlsl

@@ -0,0 +1,4 @@
+void f() {
+  int b = 0;
+  return;
+}

test/out_of_order_decls/func/type.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+fragment void f() {
+  int b = 0;
+  return;
+}
+

test/out_of_order_decls/func/type.wgsl.expected.spvasm

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %f "f"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+          %8 = OpConstantNull %int
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %b = OpVariable %_ptr_Function_int Function %8
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/func/type.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  var b : T;
+}
+
+type T = i32;

test/out_of_order_decls/func/var.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  let b = a;
+}
+
+var<private> a : i32 = 1;

test/out_of_order_decls/func/var.wgsl.expected.glsl

@@ -0,0 +1,12 @@
+#version 310 es
+precision mediump float;
+
+int a = 1;
+void f() {
+  int b = a;
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/func/var.wgsl.expected.hlsl

@@ -0,0 +1,6 @@
+static int a = 1;
+
+void f() {
+  const int b = a;
+  return;
+}

test/out_of_order_decls/func/var.wgsl.expected.msl

@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+fragment void f() {
+  thread int tint_symbol = 1;
+  int const b = tint_symbol;
+  return;
+}
+

test/out_of_order_decls/func/var.wgsl.expected.spvasm

@@ -0,0 +1,22 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %a "a"
+               OpName %f "f"
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Private_int = OpTypePointer Private %int
+          %a = OpVariable %_ptr_Private_int Private %int_1
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+          %f = OpFunction %void None %5
+          %8 = OpLabel
+          %9 = OpLoad %int %a
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/func/var.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@stage(fragment)
+fn f() {
+  let b = a;
+}
+
+var<private> a : i32 = 1;

test/out_of_order_decls/struct/alias.wgsl

@@ -0,0 +1,10 @@
+struct S {
+  m : T;
+}
+
+type T = i32;
+
+@stage(fragment)
+fn f() {
+  var v : S;
+}

test/out_of_order_decls/struct/alias.wgsl.expected.glsl

@@ -0,0 +1,15 @@
+#version 310 es
+precision mediump float;
+
+struct S {
+  int m;
+};
+
+void f() {
+  S v = S(0);
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/struct/alias.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+struct S {
+  int m;
+};
+
+void f() {
+  S v = (S)0;
+  return;
+}

test/out_of_order_decls/struct/alias.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int m;
+};
+
+fragment void f() {
+  S v = {};
+  return;
+}
+

test/out_of_order_decls/struct/alias.wgsl.expected.spvasm

@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %f "f"
+               OpName %S "S"
+               OpMemberName %S 0 "m"
+               OpName %v "v"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+%_ptr_Function_S = OpTypePointer Function %S
+          %9 = OpConstantNull %S
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_S Function %9
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/struct/alias.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+struct S {
+  m : T;
+}
+
+type T = i32;
+
+@stage(fragment)
+fn f() {
+  var v : S;
+}

test/out_of_order_decls/struct/struct.wgsl

@@ -0,0 +1,12 @@
+struct S1 {
+  m : S2;
+}
+
+struct S2 {
+  m : i32;
+}
+
+@stage(fragment)
+fn f() {
+  var v : S1;
+}

test/out_of_order_decls/struct/struct.wgsl.expected.glsl

@@ -0,0 +1,19 @@
+#version 310 es
+precision mediump float;
+
+struct S2 {
+  int m;
+};
+
+struct S1 {
+  S2 m;
+};
+
+void f() {
+  S1 v = S1(S2(0));
+}
+
+void main() {
+  f();
+  return;
+}

test/out_of_order_decls/struct/struct.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+struct S2 {
+  int m;
+};
+struct S1 {
+  S2 m;
+};
+
+void f() {
+  S1 v = (S1)0;
+  return;
+}

test/out_of_order_decls/struct/struct.wgsl.expected.msl

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S2 {
+  int m;
+};
+
+struct S1 {
+  S2 m;
+};
+
+fragment void f() {
+  S1 v = {};
+  return;
+}
+

test/out_of_order_decls/struct/struct.wgsl.expected.spvasm

@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %f "f"
+               OpExecutionMode %f OriginUpperLeft
+               OpName %f "f"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "m"
+               OpName %S2 "S2"
+               OpMemberName %S2 0 "m"
+               OpName %v "v"
+               OpMemberDecorate %S1 0 Offset 0
+               OpMemberDecorate %S2 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+         %S2 = OpTypeStruct %int
+         %S1 = OpTypeStruct %S2
+%_ptr_Function_S1 = OpTypePointer Function %S1
+         %10 = OpConstantNull %S1
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_S1 Function %10
+               OpReturn
+               OpFunctionEnd

test/out_of_order_decls/struct/struct.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+struct S1 {
+  m : S2;
+}
+
+struct S2 {
+  m : i32;
+}
+
+@stage(fragment)
+fn f() {
+  var v : S1;
+}