ProgramBuilder: Swap parameter order for Var and Const

Move the storage parameter after type.
Const() has no use for storage classes, so this parameter will be removed in the next change.
This reordering keeps Var() and Const() parameter types identical for the first two non-optional fields

Change-Id: I66669d19fa2175c4f10f615941e69efcab4c23e1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/41540
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/ast/function_test.cc b/src/ast/function_test.cc
index a74cb6a..2d65a45 100644
--- a/src/ast/function_test.cc
+++ b/src/ast/function_test.cc
@@ -31,7 +31,7 @@
 
 TEST_F(FunctionTest, Creation) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
   auto* var = params[0];
 
   auto* f = Func("func", params, ty.void_(), StatementList{},
@@ -44,7 +44,7 @@
 
 TEST_F(FunctionTest, Creation_WithSource) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f = Func(Source{Source::Location{20, 2}}, "func", params, ty.void_(),
                  StatementList{}, FunctionDecorationList{});
@@ -55,7 +55,7 @@
 
 TEST_F(FunctionTest, IsValid) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(),
                  StatementList{
@@ -67,7 +67,7 @@
 
 TEST_F(FunctionTest, IsValid_InvalidName) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f =
       Func("", params, ty.void_(), StatementList{}, FunctionDecorationList{});
@@ -76,7 +76,7 @@
 
 TEST_F(FunctionTest, IsValid_MissingReturnType) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f =
       Func("func", params, nullptr, StatementList{}, FunctionDecorationList{});
@@ -85,7 +85,7 @@
 
 TEST_F(FunctionTest, IsValid_NullParam) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
   params.push_back(nullptr);
 
   auto* f = Func("func", params, ty.void_(), StatementList{},
@@ -95,7 +95,7 @@
 
 TEST_F(FunctionTest, IsValid_InvalidParam) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, nullptr));
+  params.push_back(Var("var", nullptr, StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(), StatementList{},
                  FunctionDecorationList{});
@@ -104,7 +104,7 @@
 
 TEST_F(FunctionTest, IsValid_NullBodyStatement) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(),
                  StatementList{
@@ -118,7 +118,7 @@
 
 TEST_F(FunctionTest, IsValid_InvalidBodyStatement) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(),
                  StatementList{
@@ -162,7 +162,7 @@
 
 TEST_F(FunctionTest, ToStr_WithParams) {
   VariableList params;
-  params.push_back(Var("var", StorageClass::kNone, ty.i32()));
+  params.push_back(Var("var", ty.i32(), StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(),
                  StatementList{
@@ -192,8 +192,8 @@
 
 TEST_F(FunctionTest, TypeName_WithParams) {
   VariableList params;
-  params.push_back(Var("var1", StorageClass::kNone, ty.i32()));
-  params.push_back(Var("var2", StorageClass::kNone, ty.f32()));
+  params.push_back(Var("var1", ty.i32(), StorageClass::kNone));
+  params.push_back(Var("var2", ty.f32(), StorageClass::kNone));
 
   auto* f = Func("func", params, ty.void_(), StatementList{},
                  FunctionDecorationList{});
diff --git a/src/ast/intrinsic_texture_helper_test.cc b/src/ast/intrinsic_texture_helper_test.cc
index 27214ed..9c92015 100644
--- a/src/ast/intrinsic_texture_helper_test.cc
+++ b/src/ast/intrinsic_texture_helper_test.cc
@@ -161,27 +161,27 @@
   switch (texture_kind) {
     case ast::intrinsic::test::TextureKind::kRegular:
       return b->Global(
-          "texture", ast::StorageClass::kUniformConstant,
-          b->create<type::SampledTexture>(texture_dimension, datatype), nullptr,
-          decos);
+          "texture",
+          b->create<type::SampledTexture>(texture_dimension, datatype),
+          ast::StorageClass::kUniformConstant, nullptr, decos);
 
     case ast::intrinsic::test::TextureKind::kDepth:
-      return b->Global("texture", ast::StorageClass::kUniformConstant,
+      return b->Global("texture",
                        b->create<type::DepthTexture>(texture_dimension),
-                       nullptr, decos);
+                       ast::StorageClass::kUniformConstant, nullptr, decos);
 
     case ast::intrinsic::test::TextureKind::kMultisampled:
       return b->Global(
-          "texture", ast::StorageClass::kUniformConstant,
+          "texture",
           b->create<type::MultisampledTexture>(texture_dimension, datatype),
-          nullptr, decos);
+          ast::StorageClass::kUniformConstant, nullptr, decos);
 
     case ast::intrinsic::test::TextureKind::kStorage: {
       auto* st = b->create<type::StorageTexture>(texture_dimension,
                                                  image_format, datatype);
 
       auto* ac = b->create<type::AccessControl>(access_control, st);
-      return b->Global("texture", ast::StorageClass::kUniformConstant, ac,
+      return b->Global("texture", ac, ast::StorageClass::kUniformConstant,
                        nullptr, decos);
     }
   }
@@ -196,8 +196,8 @@
       b->create<ast::GroupDecoration>(0),
       b->create<ast::BindingDecoration>(1),
   };
-  return b->Global("sampler", ast::StorageClass::kUniformConstant,
-                   b->create<type::Sampler>(sampler_kind), nullptr, decos);
+  return b->Global("sampler", b->create<type::Sampler>(sampler_kind),
+                   ast::StorageClass::kUniformConstant, nullptr, decos);
 }
 
 std::vector<TextureOverloadCase> TextureOverloadCase::ValidCases() {
diff --git a/src/ast/module_test.cc b/src/ast/module_test.cc
index f20f4b0..7d629a9 100644
--- a/src/ast/module_test.cc
+++ b/src/ast/module_test.cc
@@ -62,7 +62,7 @@
 }
 
 TEST_F(ModuleTest, IsValid_GlobalVariable) {
-  Global("var", StorageClass::kInput, ty.f32());
+  Global("var", ty.f32(), StorageClass::kInput);
   Program program(std::move(*this));
   EXPECT_TRUE(program.AST().IsValid());
 }
@@ -74,7 +74,7 @@
 }
 
 TEST_F(ModuleTest, IsValid_Invalid_GlobalVariable) {
-  Global("var", StorageClass::kInput, nullptr);
+  Global("var", nullptr, StorageClass::kInput);
   Program program(std::move(*this));
   EXPECT_FALSE(program.AST().IsValid());
 }
diff --git a/src/ast/variable_decl_statement_test.cc b/src/ast/variable_decl_statement_test.cc
index bf18be5..15e9eb5 100644
--- a/src/ast/variable_decl_statement_test.cc
+++ b/src/ast/variable_decl_statement_test.cc
@@ -25,14 +25,14 @@
 using VariableDeclStatementTest = TestHelper;
 
 TEST_F(VariableDeclStatementTest, Creation) {
-  auto* var = Var("a", StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), StorageClass::kNone);
 
   auto* stmt = create<VariableDeclStatement>(var);
   EXPECT_EQ(stmt->variable(), var);
 }
 
 TEST_F(VariableDeclStatementTest, Creation_WithSource) {
-  auto* var = Var("a", StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), StorageClass::kNone);
 
   auto* stmt =
       create<VariableDeclStatement>(Source{Source::Location{20, 2}}, var);
@@ -42,20 +42,20 @@
 }
 
 TEST_F(VariableDeclStatementTest, IsVariableDecl) {
-  auto* var = Var("a", StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), StorageClass::kNone);
 
   auto* stmt = create<VariableDeclStatement>(var);
   EXPECT_TRUE(stmt->Is<VariableDeclStatement>());
 }
 
 TEST_F(VariableDeclStatementTest, IsValid) {
-  auto* var = Var("a", StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), StorageClass::kNone);
   auto* stmt = create<VariableDeclStatement>(var);
   EXPECT_TRUE(stmt->IsValid());
 }
 
 TEST_F(VariableDeclStatementTest, IsValid_InvalidVariable) {
-  auto* var = Var("", StorageClass::kNone, ty.f32());
+  auto* var = Var("", ty.f32(), StorageClass::kNone);
   auto* stmt = create<VariableDeclStatement>(var);
   EXPECT_FALSE(stmt->IsValid());
 }
@@ -66,7 +66,7 @@
 }
 
 TEST_F(VariableDeclStatementTest, ToStr) {
-  auto* var = Var("a", StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), StorageClass::kNone);
 
   auto* stmt =
       create<VariableDeclStatement>(Source{Source::Location{20, 2}}, var);
diff --git a/src/ast/variable_test.cc b/src/ast/variable_test.cc
index 9a9f6be..48a4c19 100644
--- a/src/ast/variable_test.cc
+++ b/src/ast/variable_test.cc
@@ -27,7 +27,7 @@
 using VariableTest = TestHelper;
 
 TEST_F(VariableTest, Creation) {
-  auto* v = Var("my_var", StorageClass::kFunction, ty.i32());
+  auto* v = Var("my_var", ty.i32(), StorageClass::kFunction);
 
   EXPECT_EQ(v->symbol(), Symbol(1));
   EXPECT_EQ(v->declared_storage_class(), StorageClass::kFunction);
@@ -41,7 +41,7 @@
 TEST_F(VariableTest, CreationWithSource) {
   auto* v = Var(
       Source{Source::Range{Source::Location{27, 4}, Source::Location{27, 5}}},
-      "i", StorageClass::kPrivate, ty.f32(), nullptr, VariableDecorationList{});
+      "i", ty.f32(), StorageClass::kPrivate, nullptr, VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
   EXPECT_EQ(v->declared_storage_class(), StorageClass::kPrivate);
@@ -55,7 +55,7 @@
 TEST_F(VariableTest, CreationEmpty) {
   auto* v = Var(
       Source{Source::Range{Source::Location{27, 4}, Source::Location{27, 7}}},
-      "a_var", StorageClass::kWorkgroup, ty.i32(), nullptr,
+      "a_var", ty.i32(), StorageClass::kWorkgroup, nullptr,
       VariableDecorationList{});
 
   EXPECT_EQ(v->symbol(), Symbol(1));
@@ -68,39 +68,39 @@
 }
 
 TEST_F(VariableTest, IsValid) {
-  auto* v = Var("my_var", StorageClass::kNone, ty.i32());
+  auto* v = Var("my_var", ty.i32(), StorageClass::kNone);
   EXPECT_TRUE(v->IsValid());
 }
 
 TEST_F(VariableTest, IsValid_WithConstructor) {
-  auto* v = Var("my_var", StorageClass::kNone, ty.i32(), Expr("ident"),
+  auto* v = Var("my_var", ty.i32(), StorageClass::kNone, Expr("ident"),
                 ast::VariableDecorationList{});
   EXPECT_TRUE(v->IsValid());
 }
 
 TEST_F(VariableTest, IsValid_MissingSymbol) {
-  auto* v = Var("", StorageClass::kNone, ty.i32());
+  auto* v = Var("", ty.i32(), StorageClass::kNone);
   EXPECT_FALSE(v->IsValid());
 }
 
 TEST_F(VariableTest, IsValid_MissingType) {
-  auto* v = Var("x", StorageClass::kNone, nullptr);
+  auto* v = Var("x", nullptr, StorageClass::kNone);
   EXPECT_FALSE(v->IsValid());
 }
 
 TEST_F(VariableTest, IsValid_MissingBoth) {
-  auto* v = Var("", StorageClass::kNone, nullptr);
+  auto* v = Var("", nullptr, StorageClass::kNone);
   EXPECT_FALSE(v->IsValid());
 }
 
 TEST_F(VariableTest, IsValid_InvalidConstructor) {
-  auto* v = Var("my_var", StorageClass::kNone, ty.i32(), Expr(""),
+  auto* v = Var("my_var", ty.i32(), StorageClass::kNone, Expr(""),
                 ast::VariableDecorationList{});
   EXPECT_FALSE(v->IsValid());
 }
 
 TEST_F(VariableTest, to_str) {
-  auto* v = Var("my_var", StorageClass::kFunction, ty.f32(), nullptr,
+  auto* v = Var("my_var", ty.f32(), StorageClass::kFunction, nullptr,
                 ast::VariableDecorationList{});
   EXPECT_EQ(str(v), R"(Variable{
   my_var
@@ -111,7 +111,7 @@
 }
 
 TEST_F(VariableTest, WithDecorations) {
-  auto* var = Var("my_var", StorageClass::kFunction, ty.i32(), nullptr,
+  auto* var = Var("my_var", ty.i32(), StorageClass::kFunction, nullptr,
                   VariableDecorationList{
                       create<LocationDecoration>(1),
                       create<BuiltinDecoration>(Builtin::kPosition),
@@ -128,7 +128,7 @@
 }
 
 TEST_F(VariableTest, ConstantId) {
-  auto* var = Var("my_var", StorageClass::kFunction, ty.i32(), nullptr,
+  auto* var = Var("my_var", ty.i32(), StorageClass::kFunction, nullptr,
                   VariableDecorationList{
                       create<ConstantIdDecoration>(1200),
                   });
@@ -137,7 +137,7 @@
 }
 
 TEST_F(VariableTest, Decorated_to_str) {
-  auto* var = Var("my_var", StorageClass::kFunction, ty.f32(), Expr("expr"),
+  auto* var = Var("my_var", ty.f32(), StorageClass::kFunction, Expr("expr"),
                   VariableDecorationList{
                       create<BindingDecoration>(2),
                       create<GroupDecoration>(1),
diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc
index 6472ebd..6fc37cb 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -109,10 +109,10 @@
       std::string in, out;
       std::tie(in, out) = inout;
 
-      Global(in, ast::StorageClass::kInput, ty.u32(), nullptr,
+      Global(in, ty.u32(), ast::StorageClass::kInput, nullptr,
              ast::VariableDecorationList{
                  create<ast::LocationDecoration>(location++)});
-      Global(out, ast::StorageClass::kOutput, ty.u32(), nullptr,
+      Global(out, ty.u32(), ast::StorageClass::kOutput, nullptr,
              ast::VariableDecorationList{
                  create<ast::LocationDecoration>(location++)});
     }
@@ -175,7 +175,7 @@
       constructor =
           create<ast::ScalarConstructorExpression>(MakeLiteral(type, val));
     }
-    GlobalConst(name, ast::StorageClass::kNone, type, constructor,
+    GlobalConst(name, type, ast::StorageClass::kNone, constructor,
                 ast::VariableDecorationList{
                     create<ast::ConstantIdDecoration>(id),
                 });
@@ -319,7 +319,7 @@
                   ast::StorageClass storage_class,
                   uint32_t group,
                   uint32_t binding) {
-    GlobalConst(name, storage_class, type, nullptr,
+    GlobalConst(name, type, storage_class, nullptr,
                 ast::VariableDecorationList{
                     create<ast::BindingDecoration>(binding),
                     create<ast::GroupDecoration>(group),
@@ -366,7 +366,7 @@
       std::string member_name = StructMemberName(member_idx, member_type);
 
       stmts.emplace_back(create<ast::VariableDeclStatement>(
-          Var("local" + member_name, ast::StorageClass::kNone, member_type)));
+          Var("local" + member_name, member_type, ast::StorageClass::kNone)));
     }
 
     for (auto member : members) {
@@ -457,14 +457,14 @@
   }
 
   void AddGlobalVariable(const std::string& name, type::Type* type) {
-    Global(name, ast::StorageClass::kUniformConstant, type);
+    Global(name, type, ast::StorageClass::kUniformConstant);
   }
 
   /// Adds a depth texture variable to the program
   /// @param name the name of the variable
   /// @param type the type to use
   void AddDepthTexture(const std::string& name, type::Type* type) {
-    Global(name, ast::StorageClass::kUniformConstant, type);
+    Global(name, type, ast::StorageClass::kUniformConstant);
   }
 
   /// Generates a function that references a specific sampler variable
@@ -486,8 +486,8 @@
 
     ast::StatementList stmts;
     stmts.emplace_back(create<ast::VariableDeclStatement>(
-        Var("sampler_result", ast::StorageClass::kFunction,
-            vec_type(base_type, 4))));
+        Var("sampler_result", vec_type(base_type, 4),
+            ast::StorageClass::kFunction)));
 
     stmts.emplace_back(create<ast::AssignmentStatement>(
         Expr("sampler_result"),
@@ -519,8 +519,8 @@
     ast::StatementList stmts;
 
     stmts.emplace_back(create<ast::VariableDeclStatement>(
-        Var("sampler_result", ast::StorageClass::kFunction,
-            vec_type(base_type, 4))));
+        Var("sampler_result", vec_type(base_type, 4),
+            ast::StorageClass::kFunction)));
 
     stmts.emplace_back(create<ast::AssignmentStatement>(
         Expr("sampler_result"), Call("textureSample", texture_name,
@@ -553,7 +553,7 @@
     ast::StatementList stmts;
 
     stmts.emplace_back(create<ast::VariableDeclStatement>(
-        Var("sampler_result", ast::StorageClass::kFunction, base_type)));
+        Var("sampler_result", base_type, ast::StorageClass::kFunction)));
     stmts.emplace_back(create<ast::AssignmentStatement>(
         Expr("sampler_result"), Call("textureSampleCompare", texture_name,
                                      sampler_name, coords_name, depth_name)));
@@ -657,7 +657,7 @@
     ast::StatementList stmts;
 
     stmts.emplace_back(create<ast::VariableDeclStatement>(
-        Var("dim", ast::StorageClass::kFunction, dim_type)));
+        Var("dim", dim_type, ast::StorageClass::kFunction)));
     stmts.emplace_back(create<ast::AssignmentStatement>(
         Expr("dim"), Call("textureDimensions", st_name)));
     stmts.emplace_back(create<ast::ReturnStatement>());
@@ -1176,10 +1176,10 @@
 }
 
 TEST_F(InspectorGetEntryPointTest, BuiltInsNotStageVariables) {
-  Global("in_var", ast::StorageClass::kInput, ty.u32(), nullptr,
+  Global("in_var", ty.u32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kPosition)});
-  Global("out_var", ast::StorageClass::kOutput, ty.u32(), nullptr,
+  Global("out_var", ty.u32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
   MakeInOutVariableBodyFunction("func", {{"in_var", "out_var"}}, {});
diff --git a/src/program_builder.h b/src/program_builder.h
index 22ccfe2..7f78bae 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -750,14 +750,14 @@
   }
 
   /// @param name the variable name
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor constructor expression
   /// @param decorations variable decorations
   /// @returns a `ast::Variable` with the given name, storage and type
   ast::Variable* Var(const std::string& name,
-                     ast::StorageClass storage,
                      type::Type* type,
+                     ast::StorageClass storage,
                      ast::Expression* constructor = nullptr,
                      ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(Symbols().Register(name), storage, type, false,
@@ -766,15 +766,15 @@
 
   /// @param source the variable source
   /// @param name the variable name
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor constructor expression
   /// @param decorations variable decorations
   /// @returns a `ast::Variable` with the given name, storage and type
   ast::Variable* Var(const Source& source,
                      const std::string& name,
-                     ast::StorageClass storage,
                      type::Type* type,
+                     ast::StorageClass storage,
                      ast::Expression* constructor = nullptr,
                      ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(source, Symbols().Register(name), storage,
@@ -782,14 +782,14 @@
   }
 
   /// @param symbol the variable symbol
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor constructor expression
   /// @param decorations variable decorations
   /// @returns a `ast::Variable` with the given symbol, storage and type
   ast::Variable* Var(Symbol symbol,
-                     ast::StorageClass storage,
                      type::Type* type,
+                     ast::StorageClass storage,
                      ast::Expression* constructor = nullptr,
                      ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(symbol, storage, type, false, constructor,
@@ -798,15 +798,15 @@
 
   /// @param source the variable source
   /// @param symbol the variable symbol
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor constructor expression
   /// @param decorations variable decorations
   /// @returns a `ast::Variable` with the given symbol, storage and type
   ast::Variable* Var(const Source& source,
                      Symbol symbol,
-                     ast::StorageClass storage,
                      type::Type* type,
+                     ast::StorageClass storage,
                      ast::Expression* constructor = nullptr,
                      ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(source, symbol, storage, type, false,
@@ -814,14 +814,14 @@
   }
 
   /// @param name the variable name
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor optional constructor expression
   /// @param decorations optional variable decorations
   /// @returns a constant `ast::Variable` with the given name, storage and type
   ast::Variable* Const(const std::string& name,
-                       ast::StorageClass storage,
                        type::Type* type,
+                       ast::StorageClass storage,
                        ast::Expression* constructor = nullptr,
                        ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(Symbols().Register(name), storage, type, true,
@@ -830,15 +830,15 @@
 
   /// @param source the variable source
   /// @param name the variable name
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor optional constructor expression
   /// @param decorations optional variable decorations
   /// @returns a constant `ast::Variable` with the given name, storage and type
   ast::Variable* Const(const Source& source,
                        const std::string& name,
-                       ast::StorageClass storage,
                        type::Type* type,
+                       ast::StorageClass storage,
                        ast::Expression* constructor = nullptr,
                        ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(source, Symbols().Register(name), storage,
@@ -846,15 +846,15 @@
   }
 
   /// @param symbol the variable symbol
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor optional constructor expression
   /// @param decorations optional variable decorations
   /// @returns a constant `ast::Variable` with the given symbol, storage and
   /// type
   ast::Variable* Const(Symbol symbol,
-                       ast::StorageClass storage,
                        type::Type* type,
+                       ast::StorageClass storage,
                        ast::Expression* constructor = nullptr,
                        ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(symbol, storage, type, true, constructor,
@@ -863,21 +863,22 @@
 
   /// @param source the variable source
   /// @param symbol the variable symbol
-  /// @param storage the variable storage class
   /// @param type the variable type
+  /// @param storage the variable storage class
   /// @param constructor optional constructor expression
   /// @param decorations optional variable decorations
   /// @returns a constant `ast::Variable` with the given symbol, storage and
   /// type
   ast::Variable* Const(const Source& source,
                        Symbol symbol,
-                       ast::StorageClass storage,
                        type::Type* type,
+                       ast::StorageClass storage,
                        ast::Expression* constructor = nullptr,
                        ast::VariableDecorationList decorations = {}) {
     return create<ast::Variable>(source, symbol, storage, type, true,
                                  constructor, decorations);
   }
+
   /// @param args the arguments to pass to Var()
   /// @returns a `ast::Variable` constructed by calling Var() with the arguments
   /// of `args`, which is automatically registered as a global variable with the
diff --git a/src/program_test.cc b/src/program_test.cc
index 6058d02..9dd64b5 100644
--- a/src/program_test.cc
+++ b/src/program_test.cc
@@ -49,7 +49,7 @@
 }
 
 TEST_F(ProgramTest, IsValid_GlobalVariable) {
-  Global("var", ast::StorageClass::kInput, ty.f32());
+  Global("var", ty.f32(), ast::StorageClass::kInput);
 
   Program program(std::move(*this));
   EXPECT_TRUE(program.IsValid());
@@ -63,7 +63,7 @@
 }
 
 TEST_F(ProgramTest, IsValid_Invalid_GlobalVariable) {
-  Global("var", ast::StorageClass::kInput, nullptr);
+  Global("var", nullptr, ast::StorageClass::kInput);
 
   Program program(std::move(*this));
   EXPECT_FALSE(program.IsValid());
diff --git a/src/scope_stack_test.cc b/src/scope_stack_test.cc
index 5e9febf..584af02 100644
--- a/src/scope_stack_test.cc
+++ b/src/scope_stack_test.cc
@@ -35,7 +35,7 @@
 }
 
 TEST_F(ScopeStackTest, Global_SetWithPointer) {
-  auto* v = Var("my_var", ast::StorageClass::kNone, ty.f32());
+  auto* v = Var("my_var", ty.f32(), ast::StorageClass::kNone);
   ScopeStack<ast::Variable*> s;
   s.set_global(v->symbol(), v);
 
diff --git a/src/transform/hlsl.cc b/src/transform/hlsl.cc
index c75eae5..da2fe0c 100644
--- a/src/transform/hlsl.cc
+++ b/src/transform/hlsl.cc
@@ -84,7 +84,7 @@
           auto* dst_init = ctx.Clone(src_init);
           // Construct the constant that holds the array
           auto* dst_var = ctx.dst->Const(
-              dst_symbol, ast::StorageClass::kFunction, dst_array_ty, dst_init);
+              dst_symbol, dst_array_ty, ast::StorageClass::kFunction, dst_init);
           // Construct the variable declaration statement
           auto* dst_var_decl =
               ctx.dst->create<ast::VariableDeclStatement>(dst_var);
diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc
index 0f3d7fb..a5a1228 100644
--- a/src/transform/spirv.cc
+++ b/src/transform/spirv.cc
@@ -68,9 +68,9 @@
         // Use `array<u32, 1>` for the new variable.
         auto* type = ctx.dst->ty.array(ctx.dst->ty.u32(), 1u);
         // Create the new variable.
-        auto* var_arr =
-            ctx.dst->Var(var->source(), var_name, var->declared_storage_class(),
-                         type, nullptr, ctx.Clone(var->decorations()));
+        auto* var_arr = ctx.dst->Var(var->source(), var_name, type,
+                                     var->declared_storage_class(), nullptr,
+                                     ctx.Clone(var->decorations()));
         // Replace the variable with the arrayed version.
         ctx.Replace(var, var_arr);
 
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc
index 71d48a3..906ded1 100644
--- a/src/type_determiner_test.cc
+++ b/src/type_determiner_test.cc
@@ -401,7 +401,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Stmt_VariableDecl) {
-  auto* var = Var("my_var", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("my_var", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
   auto* init = var->constructor();
 
@@ -416,7 +416,7 @@
 
 TEST_F(TypeDeterminerTest, Stmt_VariableDecl_ModuleScope) {
   auto* init = Expr(2);
-  Global("my_var", ast::StorageClass::kNone, ty.i32(), init,
+  Global("my_var", ty.i32(), ast::StorageClass::kNone, init,
          ast::VariableDecorationList{});
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
@@ -439,13 +439,13 @@
   ast::VariableList params;
 
   // Declare i32 "foo" inside a block
-  auto* foo_i32 = Var("foo", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* foo_i32 = Var("foo", ty.i32(), ast::StorageClass::kNone, Expr(2),
                       ast::VariableDecorationList{});
   auto* foo_i32_init = foo_i32->constructor();
   auto* foo_i32_decl = create<ast::VariableDeclStatement>(foo_i32);
 
   // Reference "foo" inside the block
-  auto* bar_i32 = Var("bar", ast::StorageClass::kNone, ty.i32(), Expr("foo"),
+  auto* bar_i32 = Var("bar", ty.i32(), ast::StorageClass::kNone, Expr("foo"),
                       ast::VariableDecorationList{});
   auto* bar_i32_init = bar_i32->constructor();
   auto* bar_i32_decl = create<ast::VariableDeclStatement>(bar_i32);
@@ -454,13 +454,13 @@
       ast::StatementList{foo_i32_decl, bar_i32_decl});
 
   // Declare f32 "foo" at function scope
-  auto* foo_f32 = Var("foo", ast::StorageClass::kNone, ty.f32(), Expr(2.f),
+  auto* foo_f32 = Var("foo", ty.f32(), ast::StorageClass::kNone, Expr(2.f),
                       ast::VariableDecorationList{});
   auto* foo_f32_init = foo_f32->constructor();
   auto* foo_f32_decl = create<ast::VariableDeclStatement>(foo_f32);
 
   // Reference "foo" at function scope
-  auto* bar_f32 = Var("bar", ast::StorageClass::kNone, ty.f32(), Expr("foo"),
+  auto* bar_f32 = Var("bar", ty.f32(), ast::StorageClass::kNone, Expr("foo"),
                       ast::VariableDecorationList{});
   auto* bar_f32_init = bar_f32->constructor();
   auto* bar_f32_decl = create<ast::VariableDeclStatement>(bar_f32);
@@ -498,7 +498,7 @@
   ast::VariableList params;
 
   // Declare i32 "foo" inside a function
-  auto* fn_i32 = Var("foo", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* fn_i32 = Var("foo", ty.i32(), ast::StorageClass::kNone, Expr(2),
                      ast::VariableDecorationList{});
   auto* fn_i32_init = fn_i32->constructor();
   auto* fn_i32_decl = create<ast::VariableDeclStatement>(fn_i32);
@@ -506,13 +506,13 @@
        ast::FunctionDecorationList{});
 
   // Declare f32 "foo" at module scope
-  auto* mod_f32 = Var("foo", ast::StorageClass::kNone, ty.f32(), Expr(2.f),
+  auto* mod_f32 = Var("foo", ty.f32(), ast::StorageClass::kNone, Expr(2.f),
                       ast::VariableDecorationList{});
   auto* mod_init = mod_f32->constructor();
   AST().AddGlobalVariable(mod_f32);
 
   // Reference "foo" in another function
-  auto* fn_f32 = Var("bar", ast::StorageClass::kNone, ty.f32(), Expr("foo"),
+  auto* fn_f32 = Var("bar", ty.f32(), ast::StorageClass::kNone, Expr("foo"),
                      ast::VariableDecorationList{});
   auto* fn_f32_init = fn_f32->constructor();
   auto* fn_f32_decl = create<ast::VariableDeclStatement>(fn_f32);
@@ -544,7 +544,7 @@
 
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Array) {
   auto* idx = Expr(2);
-  Global("my_var", ast::StorageClass::kFunction, ty.array<f32, 3>());
+  Global("my_var", ty.array<f32, 3>(), ast::StorageClass::kFunction);
 
   auto* acc = IndexAccessor("my_var", idx);
   WrapInFunction(acc);
@@ -561,7 +561,7 @@
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Alias_Array) {
   auto* aary = ty.alias("myarrty", ty.array<f32, 3>());
 
-  Global("my_var", ast::StorageClass::kFunction, aary);
+  Global("my_var", aary, ast::StorageClass::kFunction);
 
   auto* acc = IndexAccessor("my_var", 2);
   WrapInFunction(acc);
@@ -576,7 +576,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Array_Constant) {
-  GlobalConst("my_var", ast::StorageClass::kFunction, ty.array<f32, 3>());
+  GlobalConst("my_var", ty.array<f32, 3>(), ast::StorageClass::kFunction);
 
   auto* acc = IndexAccessor("my_var", 2);
   WrapInFunction(acc);
@@ -588,7 +588,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Matrix) {
-  Global("my_var", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("my_var", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* acc = IndexAccessor("my_var", 2);
   WrapInFunction(acc);
@@ -604,7 +604,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Matrix_BothDimensions) {
-  Global("my_var", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("my_var", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* acc = IndexAccessor(IndexAccessor("my_var", 2), 1);
   WrapInFunction(acc);
@@ -619,7 +619,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_ArrayAccessor_Vector) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_var", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* acc = IndexAccessor("my_var", 2);
   WrapInFunction(acc);
@@ -637,7 +637,7 @@
   auto* bitcast = create<ast::BitcastExpression>(ty.f32(), Expr("name"));
   WrapInFunction(bitcast);
 
-  Global("name", ast::StorageClass::kPrivate, ty.f32());
+  Global("name", ty.f32(), ast::StorageClass::kPrivate);
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
@@ -700,7 +700,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Cast) {
-  Global("name", ast::StorageClass::kPrivate, ty.f32());
+  Global("name", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* cast = Construct(ty.f32(), "name");
   WrapInFunction(cast);
@@ -734,7 +734,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Identifier_GlobalVariable) {
-  auto* my_var = Global("my_var", ast::StorageClass::kNone, ty.f32());
+  auto* my_var = Global("my_var", ty.f32(), ast::StorageClass::kNone);
 
   auto* ident = Expr("my_var");
   WrapInFunction(ident);
@@ -748,7 +748,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Identifier_GlobalConstant) {
-  auto* my_var = GlobalConst("my_var", ast::StorageClass::kNone, ty.f32());
+  auto* my_var = GlobalConst("my_var", ty.f32(), ast::StorageClass::kNone);
 
   auto* ident = Expr("my_var");
   WrapInFunction(ident);
@@ -763,7 +763,7 @@
 TEST_F(TypeDeterminerTest, Expr_Identifier_FunctionVariable_Const) {
   auto* my_var_a = Expr("my_var");
   auto* my_var_b = Expr("my_var");
-  auto* var = Const("my_var", ast::StorageClass::kNone, ty.f32());
+  auto* var = Const("my_var", ty.f32(), ast::StorageClass::kNone);
   auto* assign = create<ast::AssignmentStatement>(my_var_a, my_var_b);
 
   Func("my_func", ast::VariableList{}, ty.f32(),
@@ -789,7 +789,7 @@
   auto* my_var_b = Expr("my_var");
   auto* assign = create<ast::AssignmentStatement>(my_var_a, my_var_b);
 
-  auto* var = Var("my_var", ast::StorageClass::kNone, ty.f32());
+  auto* var = Var("my_var", ty.f32(), ast::StorageClass::kNone);
 
   Func("my_func", ast::VariableList{}, ty.f32(),
        ast::StatementList{
@@ -819,8 +819,8 @@
   Func("my_func", ast::VariableList{}, ty.f32(),
        ast::StatementList{
            create<ast::VariableDeclStatement>(
-               Var("my_var", ast::StorageClass::kNone,
-                   ty.pointer<f32>(ast::StorageClass::kFunction))),
+               Var("my_var", ty.pointer<f32>(ast::StorageClass::kFunction),
+                   ast::StorageClass::kNone)),
            assign,
        },
        ast::FunctionDecorationList{});
@@ -858,11 +858,11 @@
 }
 
 TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables) {
-  auto* in_var = Global("in_var", ast::StorageClass::kInput, ty.f32());
-  auto* out_var = Global("out_var", ast::StorageClass::kOutput, ty.f32());
-  auto* sb_var = Global("sb_var", ast::StorageClass::kStorage, ty.f32());
-  auto* wg_var = Global("wg_var", ast::StorageClass::kWorkgroup, ty.f32());
-  auto* priv_var = Global("priv_var", ast::StorageClass::kPrivate, ty.f32());
+  auto* in_var = Global("in_var", ty.f32(), ast::StorageClass::kInput);
+  auto* out_var = Global("out_var", ty.f32(), ast::StorageClass::kOutput);
+  auto* sb_var = Global("sb_var", ty.f32(), ast::StorageClass::kStorage);
+  auto* wg_var = Global("wg_var", ty.f32(), ast::StorageClass::kWorkgroup);
+  auto* priv_var = Global("priv_var", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* func = Func(
       "my_func", ast::VariableList{}, ty.f32(),
@@ -889,11 +889,11 @@
 }
 
 TEST_F(TypeDeterminerTest, Function_RegisterInputOutputVariables_SubFunction) {
-  auto* in_var = Global("in_var", ast::StorageClass::kInput, ty.f32());
-  auto* out_var = Global("out_var", ast::StorageClass::kOutput, ty.f32());
-  auto* sb_var = Global("sb_var", ast::StorageClass::kStorage, ty.f32());
-  auto* wg_var = Global("wg_var", ast::StorageClass::kWorkgroup, ty.f32());
-  auto* priv_var = Global("priv_var", ast::StorageClass::kPrivate, ty.f32());
+  auto* in_var = Global("in_var", ty.f32(), ast::StorageClass::kInput);
+  auto* out_var = Global("out_var", ty.f32(), ast::StorageClass::kOutput);
+  auto* sb_var = Global("sb_var", ty.f32(), ast::StorageClass::kStorage);
+  auto* wg_var = Global("wg_var", ty.f32(), ast::StorageClass::kWorkgroup);
+  auto* priv_var = Global("priv_var", ty.f32(), ast::StorageClass::kPrivate);
 
   Func("my_func", ast::VariableList{}, ty.f32(),
        ast::StatementList{
@@ -926,7 +926,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Function_NotRegisterFunctionVariable) {
-  auto* var = Var("in_var", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Var("in_var", ty.f32(), ast::StorageClass::kFunction);
 
   auto* func =
       Func("my_func", ast::VariableList{}, ty.f32(),
@@ -936,7 +936,7 @@
            },
            ast::FunctionDecorationList{});
 
-  Global("var", ast::StorageClass::kFunction, ty.f32());
+  Global("var", ty.f32(), ast::StorageClass::kFunction);
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
@@ -953,7 +953,7 @@
       ast::StructDecorationList{});
 
   auto* st = ty.struct_("S", strct);
-  Global("my_struct", ast::StorageClass::kNone, st);
+  Global("my_struct", st, ast::StorageClass::kNone);
 
   auto* mem = MemberAccessor("my_struct", "second_member");
   WrapInFunction(mem);
@@ -975,7 +975,7 @@
 
   auto* st = ty.struct_("alias", strct);
   auto* alias = ty.alias("alias", st);
-  Global("my_struct", ast::StorageClass::kNone, alias);
+  Global("my_struct", alias, ast::StorageClass::kNone);
 
   auto* mem = MemberAccessor("my_struct", "second_member");
   WrapInFunction(mem);
@@ -990,7 +990,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_MemberAccessor_VectorSwizzle) {
-  Global("my_vec", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* mem = MemberAccessor("my_vec", "xy");
   WrapInFunction(mem);
@@ -1004,7 +1004,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_MemberAccessor_VectorSwizzle_SingleElement) {
-  Global("my_vec", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* mem = MemberAccessor("my_vec", "x");
   WrapInFunction(mem);
@@ -1054,7 +1054,7 @@
       ast::StructMemberList{Member("mem", &vecB)}, ast::StructDecorationList{});
 
   auto* stA = ty.struct_("A", strctA);
-  Global("c", ast::StorageClass::kNone, stA);
+  Global("c", stA, ast::StorageClass::kNone);
 
   auto* mem = MemberAccessor(
       MemberAccessor(IndexAccessor(MemberAccessor("c", "mem"), 0), "foo"),
@@ -1076,7 +1076,7 @@
       ast::StructDecorationList{});
 
   auto* st = ty.struct_("S", strct);
-  Global("my_struct", ast::StorageClass::kNone, st);
+  Global("my_struct", st, ast::StorageClass::kNone);
 
   auto* expr = Add(MemberAccessor("my_struct", "first_member"),
                    MemberAccessor("my_struct", "second_member"));
@@ -1092,7 +1092,7 @@
 TEST_P(Expr_Binary_BitwiseTest, Scalar) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.i32());
+  Global("val", ty.i32(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1105,7 +1105,7 @@
 TEST_P(Expr_Binary_BitwiseTest, Vector) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.vec3<i32>());
+  Global("val", ty.vec3<i32>(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1132,7 +1132,7 @@
 TEST_P(Expr_Binary_LogicalTest, Scalar) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.bool_());
+  Global("val", ty.bool_(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1145,7 +1145,7 @@
 TEST_P(Expr_Binary_LogicalTest, Vector) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.vec3<bool>());
+  Global("val", ty.vec3<bool>(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1165,7 +1165,7 @@
 TEST_P(Expr_Binary_CompareTest, Scalar) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.i32());
+  Global("val", ty.i32(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1178,7 +1178,7 @@
 TEST_P(Expr_Binary_CompareTest, Vector) {
   auto op = GetParam();
 
-  Global("val", ast::StorageClass::kNone, ty.vec3<i32>());
+  Global("val", ty.vec3<i32>(), ast::StorageClass::kNone);
 
   auto* expr = create<ast::BinaryExpression>(op, Expr("val"), Expr("val"));
   WrapInFunction(expr);
@@ -1199,7 +1199,7 @@
                                          ast::BinaryOp::kGreaterThanEqual));
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Scalar_Scalar) {
-  Global("val", ast::StorageClass::kNone, ty.i32());
+  Global("val", ty.i32(), ast::StorageClass::kNone);
 
   auto* expr = Mul("val", "val");
   WrapInFunction(expr);
@@ -1210,8 +1210,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Vector_Scalar) {
-  Global("scalar", ast::StorageClass::kNone, ty.f32());
-  Global("vector", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("scalar", ty.f32(), ast::StorageClass::kNone);
+  Global("vector", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("vector", "scalar");
   WrapInFunction(expr);
@@ -1224,8 +1224,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Scalar_Vector) {
-  Global("scalar", ast::StorageClass::kNone, ty.f32());
-  Global("vector", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("scalar", ty.f32(), ast::StorageClass::kNone);
+  Global("vector", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("scalar", "vector");
   WrapInFunction(expr);
@@ -1238,7 +1238,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Vector_Vector) {
-  Global("vector", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("vector", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("vector", "vector");
   WrapInFunction(expr);
@@ -1251,8 +1251,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Matrix_Scalar) {
-  Global("scalar", ast::StorageClass::kNone, ty.f32());
-  Global("matrix", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("scalar", ty.f32(), ast::StorageClass::kNone);
+  Global("matrix", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("matrix", "scalar");
   WrapInFunction(expr);
@@ -1268,8 +1268,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Scalar_Matrix) {
-  Global("scalar", ast::StorageClass::kNone, ty.f32());
-  Global("matrix", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("scalar", ty.f32(), ast::StorageClass::kNone);
+  Global("matrix", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("scalar", "matrix");
   WrapInFunction(expr);
@@ -1285,8 +1285,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Matrix_Vector) {
-  Global("vector", ast::StorageClass::kNone, ty.vec3<f32>());
-  Global("matrix", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("vector", ty.vec3<f32>(), ast::StorageClass::kNone);
+  Global("matrix", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("matrix", "vector");
   WrapInFunction(expr);
@@ -1299,8 +1299,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Vector_Matrix) {
-  Global("vector", ast::StorageClass::kNone, ty.vec3<f32>());
-  Global("matrix", ast::StorageClass::kNone, ty.mat2x3<f32>());
+  Global("vector", ty.vec3<f32>(), ast::StorageClass::kNone);
+  Global("matrix", ty.mat2x3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("vector", "matrix");
   WrapInFunction(expr);
@@ -1313,8 +1313,8 @@
 }
 
 TEST_F(TypeDeterminerTest, Expr_Binary_Multiply_Matrix_Matrix) {
-  Global("mat3x4", ast::StorageClass::kNone, ty.mat3x4<f32>());
-  Global("mat4x3", ast::StorageClass::kNone, ty.mat4x3<f32>());
+  Global("mat3x4", ty.mat3x4<f32>(), ast::StorageClass::kNone);
+  Global("mat4x3", ty.mat4x3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Mul("mat3x4", "mat4x3");
   WrapInFunction(expr);
@@ -1333,7 +1333,7 @@
 TEST_P(IntrinsicDerivativeTest, Scalar) {
   auto name = GetParam();
 
-  Global("ident", ast::StorageClass::kNone, ty.f32());
+  Global("ident", ty.f32(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "ident");
   WrapInFunction(expr);
@@ -1346,7 +1346,7 @@
 
 TEST_P(IntrinsicDerivativeTest, Vector) {
   auto name = GetParam();
-  Global("ident", ast::StorageClass::kNone, ty.vec4<f32>());
+  Global("ident", ty.vec4<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "ident");
   WrapInFunction(expr);
@@ -1390,7 +1390,7 @@
 TEST_P(Intrinsic, Test) {
   auto name = GetParam();
 
-  Global("my_var", ast::StorageClass::kNone, ty.vec3<bool>());
+  Global("my_var", ty.vec3<bool>(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "my_var");
   WrapInFunction(expr);
@@ -1408,7 +1408,7 @@
 TEST_P(Intrinsic_FloatMethod, Vector) {
   auto name = GetParam();
 
-  Global("my_var", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_var", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "my_var");
   WrapInFunction(expr);
@@ -1424,7 +1424,7 @@
 TEST_P(Intrinsic_FloatMethod, Scalar) {
   auto name = GetParam();
 
-  Global("my_var", ast::StorageClass::kNone, ty.f32());
+  Global("my_var", ty.f32(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "my_var");
   WrapInFunction(expr);
@@ -1438,7 +1438,7 @@
 TEST_P(Intrinsic_FloatMethod, MissingParam) {
   auto name = GetParam();
 
-  Global("my_var", ast::StorageClass::kNone, ty.f32());
+  Global("my_var", ty.f32(), ast::StorageClass::kNone);
 
   auto* expr = Call(name);
   WrapInFunction(expr);
@@ -1455,7 +1455,7 @@
 TEST_P(Intrinsic_FloatMethod, TooManyParams) {
   auto name = GetParam();
 
-  Global("my_var", ast::StorageClass::kNone, ty.f32());
+  Global("my_var", ty.f32(), ast::StorageClass::kNone);
 
   auto* expr = Call(name, "my_var", 1.23f);
   WrapInFunction(expr);
@@ -1524,7 +1524,7 @@
   void add_call_param(std::string name,
                       type::Type* type,
                       ast::ExpressionList* call_params) {
-    Global(name, ast::StorageClass::kNone, type);
+    Global(name, type, ast::StorageClass::kNone);
     call_params->push_back(Expr(name));
   }
   type::Type* subtype(Texture type) {
@@ -1654,7 +1654,7 @@
                     TextureTestParams{type::TextureDimension::k3d}));
 
 TEST_F(TypeDeterminerTest, Intrinsic_Dot_Vec2) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec2<f32>());
+  Global("my_var", ty.vec2<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -1666,7 +1666,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Dot_Vec3) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_var", ty.vec3<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -1678,7 +1678,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Dot_Vec4) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec4<f32>());
+  Global("my_var", ty.vec4<f32>(), ast::StorageClass::kNone);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -1704,7 +1704,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Dot_Error_VectorInt) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec4<i32>());
+  Global("my_var", ty.vec4<i32>(), ast::StorageClass::kNone);
 
   auto* expr = Call("dot", "my_var", "my_var");
   WrapInFunction(expr);
@@ -1720,9 +1720,9 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Select) {
-  Global("my_var", ast::StorageClass::kNone, ty.vec3<f32>());
+  Global("my_var", ty.vec3<f32>(), ast::StorageClass::kNone);
 
-  Global("bool_var", ast::StorageClass::kNone, ty.vec3<bool>());
+  Global("bool_var", ty.vec3<bool>(), ast::StorageClass::kNone);
 
   auto* expr = Call("select", "my_var", "my_var", "bool_var");
   WrapInFunction(expr);
@@ -1816,7 +1816,7 @@
 TEST_P(UnaryOpExpressionTest, Expr_UnaryOp) {
   auto op = GetParam();
 
-  Global("ident", ast::StorageClass::kNone, ty.vec4<f32>());
+  Global("ident", ty.vec4<f32>(), ast::StorageClass::kNone);
   auto* der = create<ast::UnaryOpExpression>(op, Expr("ident"));
   WrapInFunction(der);
 
@@ -1833,7 +1833,7 @@
                                          ast::UnaryOp::kNot));
 
 TEST_F(TypeDeterminerTest, StorageClass_SetsIfMissing) {
-  auto* var = Var("var", ast::StorageClass::kNone, ty.i32());
+  auto* var = Var("var", ty.i32(), ast::StorageClass::kNone);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   Func("func", ast::VariableList{}, ty.i32(), ast::StatementList{stmt},
@@ -1845,7 +1845,7 @@
 }
 
 TEST_F(TypeDeterminerTest, StorageClass_DoesNotSetOnConst) {
-  auto* var = Const("var", ast::StorageClass::kNone, ty.i32());
+  auto* var = Const("var", ty.i32(), ast::StorageClass::kNone);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   Func("func", ast::VariableList{}, ty.i32(), ast::StatementList{stmt},
        ast::FunctionDecorationList{});
@@ -1856,7 +1856,7 @@
 }
 
 TEST_F(TypeDeterminerTest, StorageClass_NonFunctionClassError) {
-  auto* var = Var("var", ast::StorageClass::kWorkgroup, ty.i32());
+  auto* var = Var("var", ty.i32(), ast::StorageClass::kWorkgroup);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   Func("func", ast::VariableList{}, ty.i32(), ast::StatementList{stmt},
@@ -2148,7 +2148,7 @@
                     IntrinsicData{"trunc", IntrinsicType::kTrunc}));
 
 TEST_F(IntrinsicDataTest, ArrayLength_Vector) {
-  Global("arr", ast::StorageClass::kNone, ty.array<int>());
+  Global("arr", ty.array<int>(), ast::StorageClass::kNone);
   auto* call = Call("arrayLength", "arr");
   WrapInFunction(call);
 
@@ -2159,7 +2159,7 @@
 }
 
 TEST_F(IntrinsicDataTest, ArrayLength_Error_ArraySized) {
-  Global("arr", ast::StorageClass::kNone, ty.array<int, 4>());
+  Global("arr", ty.array<int, 4>(), ast::StorageClass::kNone);
   auto* call = Call("arrayLength", "arr");
   WrapInFunction(call);
 
@@ -2195,7 +2195,7 @@
 }
 
 TEST_F(IntrinsicDataTest, FrexpScalar) {
-  Global("exp", ast::StorageClass::kWorkgroup, ty.i32());
+  Global("exp", ty.i32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", 1.0f, "exp");
   WrapInFunction(call);
 
@@ -2206,7 +2206,7 @@
 }
 
 TEST_F(IntrinsicDataTest, FrexpVector) {
-  Global("exp", ast::StorageClass::kWorkgroup, ty.vec3<i32>());
+  Global("exp", ty.vec3<i32>(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", vec3<f32>(1.0f, 2.0f, 3.0f), "exp");
   WrapInFunction(call);
 
@@ -2218,7 +2218,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Frexp_Error_FirstParamInt) {
-  Global("exp", ast::StorageClass::kWorkgroup, ty.i32());
+  Global("exp", ty.i32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", 1, "exp");
   WrapInFunction(call);
 
@@ -2233,7 +2233,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Frexp_Error_SecondParamFloatPtr) {
-  Global("exp", ast::StorageClass::kWorkgroup, ty.f32());
+  Global("exp", ty.f32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", 1.0f, "exp");
   WrapInFunction(call);
 
@@ -2262,7 +2262,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Frexp_Error_VectorSizesDontMatch) {
-  Global("exp", ast::StorageClass::kWorkgroup, ty.vec4<i32>());
+  Global("exp", ty.vec4<i32>(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", vec2<f32>(1.0f, 2.0f), "exp");
   WrapInFunction(call);
 
@@ -2278,7 +2278,7 @@
 }
 
 TEST_F(IntrinsicDataTest, ModfScalar) {
-  Global("whole", ast::StorageClass::kWorkgroup, ty.f32());
+  Global("whole", ty.f32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", 1.0f, "whole");
   WrapInFunction(call);
 
@@ -2289,7 +2289,7 @@
 }
 
 TEST_F(IntrinsicDataTest, ModfVector) {
-  Global("whole", ast::StorageClass::kWorkgroup, ty.vec3<f32>());
+  Global("whole", ty.vec3<f32>(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", vec3<f32>(1.0f, 2.0f, 3.0f), "whole");
   WrapInFunction(call);
 
@@ -2301,7 +2301,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Modf_Error_FirstParamInt) {
-  Global("whole", ast::StorageClass::kWorkgroup, ty.f32());
+  Global("whole", ty.f32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", 1, "whole");
   WrapInFunction(call);
 
@@ -2315,7 +2315,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Modf_Error_SecondParamIntPtr) {
-  Global("whole", ast::StorageClass::kWorkgroup, ty.i32());
+  Global("whole", ty.i32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", 1.0f, "whole");
   WrapInFunction(call);
 
@@ -2342,7 +2342,7 @@
 }
 
 TEST_F(IntrinsicDataTest, Modf_Error_VectorSizesDontMatch) {
-  Global("whole", ast::StorageClass::kWorkgroup, ty.vec4<f32>());
+  Global("whole", ty.vec4<f32>(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", vec2<f32>(1.0f, 2.0f), "whole");
   WrapInFunction(call);
 
@@ -2993,7 +2993,7 @@
                     IntrinsicData{"max", IntrinsicType::kMax}));
 
 TEST_F(TypeDeterminerTest, Intrinsic_Determinant_2x2) {
-  Global("var", ast::StorageClass::kFunction, ty.mat2x2<f32>());
+  Global("var", ty.mat2x2<f32>(), ast::StorageClass::kFunction);
 
   auto* call = Call("determinant", "var");
   WrapInFunction(call);
@@ -3005,7 +3005,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Determinant_3x3) {
-  Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("var", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* call = Call("determinant", "var");
   WrapInFunction(call);
@@ -3017,7 +3017,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Determinant_4x4) {
-  Global("var", ast::StorageClass::kFunction, ty.mat4x4<f32>());
+  Global("var", ty.mat4x4<f32>(), ast::StorageClass::kFunction);
 
   auto* call = Call("determinant", "var");
   WrapInFunction(call);
@@ -3029,7 +3029,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Determinant_NotSquare) {
-  Global("var", ast::StorageClass::kFunction, ty.mat2x3<f32>());
+  Global("var", ty.mat2x3<f32>(), ast::StorageClass::kFunction);
 
   auto* call = Call("determinant", "var");
   WrapInFunction(call);
@@ -3043,7 +3043,7 @@
 }
 
 TEST_F(TypeDeterminerTest, Intrinsic_Determinant_NotMatrix) {
-  Global("var", ast::StorageClass::kFunction, ty.f32());
+  Global("var", ty.f32(), ast::StorageClass::kFunction);
 
   auto* call = Call("determinant", "var");
   WrapInFunction(call);
@@ -3105,11 +3105,11 @@
                create<ast::StageDecoration>(ast::PipelineStage::kVertex),
            });
 
-  Global("first", ast::StorageClass::kPrivate, ty.f32());
-  Global("second", ast::StorageClass::kPrivate, ty.f32());
-  Global("call_a", ast::StorageClass::kPrivate, ty.f32());
-  Global("call_b", ast::StorageClass::kPrivate, ty.f32());
-  Global("call_c", ast::StorageClass::kPrivate, ty.f32());
+  Global("first", ty.f32(), ast::StorageClass::kPrivate);
+  Global("second", ty.f32(), ast::StorageClass::kPrivate);
+  Global("call_a", ty.f32(), ast::StorageClass::kPrivate);
+  Global("call_b", ty.f32(), ast::StorageClass::kPrivate);
+  Global("call_c", ty.f32(), ast::StorageClass::kPrivate);
 
   ASSERT_TRUE(td()->Determine()) << td()->error();
 
diff --git a/src/validator/validator_builtins_test.cc b/src/validator/validator_builtins_test.cc
index f0b7924..dee228b 100644
--- a/src/validator/validator_builtins_test.cc
+++ b/src/validator/validator_builtins_test.cc
@@ -144,7 +144,7 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Frexp_Scalar) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.i32());
+  auto* a = Var("a", ty.i32(), ast::StorageClass::kWorkgroup);
   RegisterVariable(a);
   auto* builtin = Call("frexp", 1.0f, Expr("a"));
   WrapInFunction(builtin);
@@ -157,11 +157,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Frexp_Vec2) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec2<int>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec2<int>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.i32(), 2),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("frexp", vec2<float>(1.0f, 1.0f), Expr("b"));
@@ -176,11 +176,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Frexp_Vec3) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec3<int>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec3<int>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.i32(), 3),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("frexp", vec3<float>(1.0f, 1.0f, 1.0f), Expr("b"));
@@ -194,11 +194,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Frexp_Vec4) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec4<int>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec4<int>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.i32(), 4),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("frexp", vec4<float>(1.0f, 1.0f, 1.0f, 1.0f), Expr("b"));
@@ -212,10 +212,9 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Modf_Scalar) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.f32());
-  auto* b =
-      Const("b", ast::StorageClass::kWorkgroup,
-            ty.pointer<float>(ast::StorageClass::kWorkgroup), Expr("a"), {});
+  auto* a = Var("a", ty.f32(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b", ty.pointer<float>(ast::StorageClass::kWorkgroup),
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("modf", 1.0f, Expr("b"));
@@ -229,11 +228,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Modf_Vec2) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec2<float>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec2<float>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.f32(), 2),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("modf", vec2<float>(1.0f, 1.0f), Expr("b"));
@@ -247,11 +246,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Modf_Vec3) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec3<float>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec3<float>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.f32(), 3),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("modf", vec3<float>(1.0f, 1.0f, 1.0f), Expr("b"));
@@ -265,11 +264,11 @@
 }
 
 TEST_F(ValidatorBuiltinsTest, Modf_Vec4) {
-  auto* a = Var("a", ast::StorageClass::kWorkgroup, ty.vec4<float>());
-  auto* b = Const("b", ast::StorageClass::kWorkgroup,
+  auto* a = Var("a", ty.vec4<float>(), ast::StorageClass::kWorkgroup);
+  auto* b = Const("b",
                   create<type::Pointer>(create<type::Vector>(ty.f32(), 4),
                                         ast::StorageClass::kWorkgroup),
-                  Expr("a"), {});
+                  ast::StorageClass::kWorkgroup, Expr("a"), {});
   RegisterVariable(a);
   RegisterVariable(b);
   auto* builtin = Call("modf", vec4<float>(1.0f, 1.0f, 1.0f, 1.0f), Expr("b"));
diff --git a/src/validator/validator_control_block_test.cc b/src/validator/validator_control_block_test.cc
index 6a76b5d..5852ad8 100644
--- a/src/validator/validator_control_block_test.cc
+++ b/src/validator/validator_control_block_test.cc
@@ -42,7 +42,7 @@
   // switch (a) {
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(3.14f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(3.14f),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList body;
@@ -71,7 +71,7 @@
   // switch (a) {
   //   case 1: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseSelectorList csl;
@@ -104,7 +104,7 @@
   //   case 1: {}
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList switch_body;
@@ -146,7 +146,7 @@
   //   case 1: {}
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList switch_body;
@@ -181,7 +181,7 @@
   //   case -1: {}
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.u32(), Expr(2u),
+  auto* var = Var("a", ty.u32(), ast::StorageClass::kNone, Expr(2u),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList switch_body;
@@ -216,7 +216,7 @@
   //   case 2, 2: {}
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.u32(), Expr(3u),
+  auto* var = Var("a", ty.u32(), ast::StorageClass::kNone, Expr(3u),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList switch_body;
@@ -257,7 +257,7 @@
   //   case 0,1,2,10: {}
   //   default: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseStatementList switch_body;
@@ -298,7 +298,7 @@
   // switch (a) {
   //   default: { fallthrough; }
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseSelectorList default_csl;
@@ -330,7 +330,7 @@
   //   default: {}
   //   case 5: {}
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   ast::CaseSelectorList default_csl;
@@ -362,7 +362,7 @@
   // }
 
   auto* my_int = ty.alias("MyInt", ty.u32());
-  auto* var = Var("a", ast::StorageClass::kNone, my_int, Expr(2u),
+  auto* var = Var("a", my_int, ast::StorageClass::kNone, Expr(2u),
                   ast::VariableDecorationList{});
 
   ast::CaseSelectorList default_csl;
diff --git a/src/validator/validator_function_test.cc b/src/validator/validator_function_test.cc
index 13f6e59..74e4bb2 100644
--- a/src/validator/validator_function_test.cc
+++ b/src/validator/validator_function_test.cc
@@ -38,7 +38,7 @@
 TEST_F(ValidateFunctionTest, VoidFunctionEndWithoutReturnStatement_Pass) {
   // [[stage(vertex)]]
   // fn func -> void { var a:i32 = 2; }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   Func(Source{Source::Location{12, 34}}, "func", ast::VariableList{},
@@ -74,7 +74,7 @@
 TEST_F(ValidateFunctionTest, FunctionEndWithoutReturnStatement_Fail) {
   // fn func -> int { var a:i32 = 2; }
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   Func(Source{Source::Location{12, 34}}, "func", ast::VariableList{}, ty.i32(),
@@ -201,7 +201,7 @@
   ast::ExpressionList call_params;
   auto* call_expr = create<ast::CallExpression>(
       Source{Source::Location{12, 34}}, Expr("func"), call_params);
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), call_expr,
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, call_expr,
                   ast::VariableDecorationList{});
 
   Func("func", ast::VariableList{}, ty.i32(),
@@ -242,7 +242,7 @@
   // fn vtx_func(a : i32) -> void { return; }
 
   Func(Source{Source::Location{12, 34}}, "vtx_func",
-       ast::VariableList{Var("a", ast::StorageClass::kNone, ty.i32(), nullptr,
+       ast::VariableList{Var("a", ty.i32(), ast::StorageClass::kNone, nullptr,
                              ast::VariableDecorationList{})},
        ty.void_(),
        ast::StatementList{
diff --git a/src/validator/validator_test.cc b/src/validator/validator_test.cc
index c41bc16..6ca6e6e 100644
--- a/src/validator/validator_test.cc
+++ b/src/validator/validator_test.cc
@@ -64,7 +64,7 @@
   // var my_var : i32 = 2;
   // 1 = my_var;
 
-  auto* var = Var("my_var", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("my_var", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
   RegisterVariable(var);
 
@@ -122,7 +122,7 @@
 TEST_F(ValidatorTest, AssignCompatibleTypes_Pass) {
   // var a :i32 = 2;
   // a = 2
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
   RegisterVariable(var);
 
@@ -146,7 +146,7 @@
   // var a :myint = 2;
   // a = 2
   auto* myint = ty.alias("myint", ty.i32());
-  auto* var = Var("a", ast::StorageClass::kNone, myint, Expr(2),
+  auto* var = Var("a", myint, ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
   RegisterVariable(var);
 
@@ -169,9 +169,9 @@
   // var a :i32 = 2;
   // var b :i32 = 3;
   // a = b;
-  auto* var_a = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var_a = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                     ast::VariableDecorationList{});
-  auto* var_b = Var("b", ast::StorageClass::kNone, ty.i32(), Expr(3),
+  auto* var_b = Var("b", ty.i32(), ast::StorageClass::kNone, Expr(3),
                     ast::VariableDecorationList{});
   RegisterVariable(var_a);
   RegisterVariable(var_b);
@@ -196,8 +196,8 @@
   // const b : ptr<function,i32> = a;
   // b = 2;
   const auto func = ast::StorageClass::kFunction;
-  auto* var_a = Var("a", func, ty.i32(), Expr(2), {});
-  auto* var_b = Const("b", ast::StorageClass::kNone, ty.pointer<int>(func),
+  auto* var_a = Var("a", ty.i32(), func, Expr(2), {});
+  auto* var_b = Const("b", ty.pointer<int>(func), ast::StorageClass::kNone,
                       Expr("a"), {});
   RegisterVariable(var_a);
   RegisterVariable(var_b);
@@ -223,7 +223,7 @@
   //  a = 2.3;
   // }
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
   RegisterVariable(var);
 
@@ -252,8 +252,8 @@
   // const b : ptr<function,f32> = a;
   // b = 2;
   const auto priv = ast::StorageClass::kFunction;
-  auto* var_a = Var("a", priv, ty.f32(), Expr(2), {});
-  auto* var_b = Const("b", ast::StorageClass::kNone, ty.pointer<float>(priv),
+  auto* var_a = Var("a", ty.f32(), priv, Expr(2), {});
+  auto* var_b = Const("b", ty.pointer<float>(priv), ast::StorageClass::kNone,
                       Expr("a"), {});
   RegisterVariable(var_a);
   RegisterVariable(var_b);
@@ -281,7 +281,7 @@
   //  var a :i32 = 2;
   //  a = 2
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   auto* lhs = Expr("a");
@@ -308,7 +308,7 @@
   //  a = 2.3;
   // }
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   auto* lhs = Expr("a");
@@ -342,7 +342,7 @@
   //  }
   // }
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
   auto* lhs = Expr("a");
@@ -375,9 +375,9 @@
 
 TEST_F(ValidatorTest, GlobalVariableWithStorageClass_Pass) {
   // var<in> gloabl_var: f32;
-  auto* var = Global(Source{Source::Location{12, 34}}, "global_var",
-                     ast::StorageClass::kInput, ty.f32(), nullptr,
-                     ast::VariableDecorationList{});
+  auto* var =
+      Global(Source{Source::Location{12, 34}}, "global_var", ty.f32(),
+             ast::StorageClass::kInput, nullptr, ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
 
@@ -386,9 +386,8 @@
 
 TEST_F(ValidatorTest, GlobalVariableNoStorageClass_Fail) {
   // var gloabl_var: f32;
-  Global(Source{Source::Location{12, 34}}, "global_var",
-         ast::StorageClass::kNone, ty.f32(), nullptr,
-         ast::VariableDecorationList{});
+  Global(Source{Source::Location{12, 34}}, "global_var", ty.f32(),
+         ast::StorageClass::kNone, nullptr, ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
 
@@ -399,8 +398,8 @@
 
 TEST_F(ValidatorTest, GlobalConstantWithStorageClass_Fail) {
   // const<in> gloabl_var: f32;
-  GlobalConst(Source{Source::Location{12, 34}}, "global_var",
-              ast::StorageClass::kInput, ty.f32(), nullptr,
+  GlobalConst(Source{Source::Location{12, 34}}, "global_var", ty.f32(),
+              ast::StorageClass::kInput, nullptr,
               ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
@@ -413,9 +412,8 @@
 
 TEST_F(ValidatorTest, GlobalConstNoStorageClass_Pass) {
   // const gloabl_var: f32;
-  GlobalConst(Source{Source::Location{12, 34}}, "global_var",
-              ast::StorageClass::kNone, ty.f32(), nullptr,
-              ast::VariableDecorationList{});
+  GlobalConst(Source{Source::Location{12, 34}}, "global_var", ty.f32(),
+              ast::StorageClass::kNone, nullptr, ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
 
@@ -439,7 +437,7 @@
        ast::FunctionDecorationList{
            create<ast::StageDecoration>(ast::PipelineStage::kVertex)});
 
-  Global("global_var", ast::StorageClass::kPrivate, ty.f32(), Expr(2.1f),
+  Global("global_var", ty.f32(), ast::StorageClass::kPrivate, Expr(2.1f),
          ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
@@ -455,7 +453,7 @@
   //   return;
   // }
 
-  Global("global_var", ast::StorageClass::kPrivate, ty.f32(), Expr(2.1f),
+  Global("global_var", ty.f32(), ast::StorageClass::kPrivate, Expr(2.1f),
          ast::VariableDecorationList{});
 
   Func("my_func", ast::VariableList{}, ty.void_(),
@@ -478,7 +476,7 @@
   //   if (true) { var a : f32 = 2.0; }
   //   a = 3.14;
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   auto* cond = Expr(true);
@@ -512,7 +510,7 @@
   //   var a : f32 = 2.0;
   //   if (true) { a = 3.14; }
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   SetSource(Source{Source::Location{12, 34}});
@@ -545,7 +543,7 @@
   //  { var a : f32 = 2.0; }
   //  { a = 3.14; }
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
   auto* first_body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(var),
@@ -578,11 +576,11 @@
 TEST_F(ValidatorTest, GlobalVariableUnique_Pass) {
   // var global_var0 : f32 = 0.1;
   // var global_var1 : i32 = 0;
-  auto* var0 = Global("global_var0", ast::StorageClass::kPrivate, ty.f32(),
+  auto* var0 = Global("global_var0", ty.f32(), ast::StorageClass::kPrivate,
                       Expr(0.1f), ast::VariableDecorationList{});
 
-  auto* var1 = Global(Source{Source::Location{12, 34}}, "global_var1",
-                      ast::StorageClass::kPrivate, ty.f32(), Expr(0),
+  auto* var1 = Global(Source{Source::Location{12, 34}}, "global_var1", ty.f32(),
+                      ast::StorageClass::kPrivate, Expr(0),
                       ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
@@ -594,12 +592,11 @@
 TEST_F(ValidatorTest, GlobalVariableNotUnique_Fail) {
   // var global_var : f32 = 0.1;
   // var global_var : i32 = 0;
-  Global("global_var", ast::StorageClass::kPrivate, ty.f32(), Expr(0.1f),
+  Global("global_var", ty.f32(), ast::StorageClass::kPrivate, Expr(0.1f),
          ast::VariableDecorationList{});
 
-  Global(Source{Source::Location{12, 34}}, "global_var",
-         ast::StorageClass::kPrivate, ty.i32(), Expr(0),
-         ast::VariableDecorationList{});
+  Global(Source{Source::Location{12, 34}}, "global_var", ty.i32(),
+         ast::StorageClass::kPrivate, Expr(0), ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
 
@@ -613,7 +610,7 @@
   //  const a :i32 = 2;
   //  a = 2
   // }
-  auto* var = Const("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Const("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                     ast::VariableDecorationList{});
 
   auto* lhs = Expr("a");
@@ -642,7 +639,7 @@
   // }
   // var a: f32 = 2.1;
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   Func("my_func", ast::VariableList{}, ty.void_(),
@@ -652,7 +649,7 @@
        ast::FunctionDecorationList{
            create<ast::StageDecoration>(ast::PipelineStage::kVertex)});
 
-  Global("a", ast::StorageClass::kPrivate, ty.f32(), Expr(2.1f),
+  Global("a", ty.f32(), ast::StorageClass::kPrivate, Expr(2.1f),
          ast::VariableDecorationList{});
 
   ValidatorImpl& v = Build();
@@ -667,10 +664,10 @@
   //   return 0;
   // }
 
-  Global("a", ast::StorageClass::kPrivate, ty.f32(), Expr(2.1f),
+  Global("a", ty.f32(), ast::StorageClass::kPrivate, Expr(2.1f),
          ast::VariableDecorationList{});
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   Func("my_func", ast::VariableList{}, ty.void_(),
@@ -691,10 +688,10 @@
   //  var a :i32 = 2;
   //  var a :f21 = 2.0;
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, Expr(2),
                   ast::VariableDecorationList{});
 
-  auto* var_a_float = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(0.1f),
+  auto* var_a_float = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(0.1f),
                           ast::VariableDecorationList{});
 
   Func("my_func", ast::VariableList{}, ty.void_(),
@@ -716,7 +713,7 @@
   // if (true) { var a : f32 = 2.0; }
   // var a : f32 = 3.14;
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   auto* cond = Expr(true);
@@ -724,7 +721,7 @@
       create<ast::VariableDeclStatement>(var),
   });
 
-  auto* var_a_float = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(3.1f),
+  auto* var_a_float = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(3.1f),
                           ast::VariableDecorationList{});
 
   auto* outer_body = create<ast::BlockStatement>(ast::StatementList{
@@ -747,10 +744,10 @@
   // var a : f32 = 3.14;
   // if (true) { var a : f32 = 2.0; }
   // }
-  auto* var_a_float = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(3.1f),
+  auto* var_a_float = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(3.1f),
                           ast::VariableDecorationList{});
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                   ast::VariableDecorationList{});
 
   auto* cond = Expr(true);
@@ -776,13 +773,13 @@
   //  { var a : f32; }
   //  var a : f32;
   // }
-  auto* var_inner = Var("a", ast::StorageClass::kNone, ty.f32());
+  auto* var_inner = Var("a", ty.f32(), ast::StorageClass::kNone);
   auto* inner = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(Source{Source::Location{12, 34}},
                                          var_inner),
   });
 
-  auto* var_outer = Var("a", ast::StorageClass::kNone, ty.f32());
+  auto* var_outer = Var("a", ty.f32(), ast::StorageClass::kNone);
   auto* outer_body = create<ast::BlockStatement>(ast::StatementList{
       inner,
       create<ast::VariableDeclStatement>(var_outer),
@@ -800,13 +797,13 @@
   //  var a : f32;
   //  { var a : f32; }
   // }
-  auto* var_inner = Var("a", ast::StorageClass::kNone, ty.f32());
+  auto* var_inner = Var("a", ty.f32(), ast::StorageClass::kNone);
   auto* inner = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(Source{Source::Location{12, 34}},
                                          var_inner),
   });
 
-  auto* var_outer = Var("a", ast::StorageClass::kNone, ty.f32());
+  auto* var_outer = Var("a", ty.f32(), ast::StorageClass::kNone);
   auto* outer_body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(var_outer),
       inner,
@@ -823,10 +820,10 @@
 TEST_F(ValidatorTest, RedeclaredIdentifierDifferentFunctions_Pass) {
   // func0 { var a : f32 = 2.0; return; }
   // func1 { var a : f32 = 3.0; return; }
-  auto* var0 = Var("a", ast::StorageClass::kNone, ty.f32(), Expr(2.0f),
+  auto* var0 = Var("a", ty.f32(), ast::StorageClass::kNone, Expr(2.0f),
                    ast::VariableDecorationList{});
 
-  auto* var1 = Var("a", ast::StorageClass::kNone, ty.void_(), Expr(1.0f),
+  auto* var1 = Var("a", ty.void_(), ast::StorageClass::kNone, Expr(1.0f),
                    ast::VariableDecorationList{});
 
   Func("func0", ast::VariableList{}, ty.void_(),
@@ -857,7 +854,7 @@
   // var a :i32;
   // a = 2;
   // }
-  auto* var = Var("a", ast::StorageClass::kNone, ty.i32(), nullptr,
+  auto* var = Var("a", ty.i32(), ast::StorageClass::kNone, nullptr,
                   ast::VariableDecorationList{});
   auto* lhs = Expr("a");
   auto* rhs = Expr(2);
diff --git a/src/validator/validator_type_test.cc b/src/validator/validator_type_test.cc
index 1df5059..18fa10d 100644
--- a/src/validator/validator_type_test.cc
+++ b/src/validator/validator_type_test.cc
@@ -162,7 +162,7 @@
   /// [[stage(vertex)]]
   // fn func -> void { var a : array<i32>; }
 
-  auto* var = Var("a", ast::StorageClass::kNone, ty.array<i32>());
+  auto* var = Var("a", ty.array<i32>(), ast::StorageClass::kNone);
 
   Func("func", ast::VariableList{}, ty.void_(),
        ast::StatementList{
@@ -186,8 +186,8 @@
   // [[stage(vertex)]] fn main() {}
 
   auto* param =
-      Var(Source{Source::Location{12, 34}}, "a", ast::StorageClass::kNone,
-          ty.array<i32>(), nullptr, ast::VariableDecorationList{});
+      Var(Source{Source::Location{12, 34}}, "a", ty.array<i32>(),
+          ast::StorageClass::kNone, nullptr, ast::VariableDecorationList{});
 
   Func("func", ast::VariableList{param}, ty.void_(),
        ast::StatementList{
diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc
index bbe1eca..425d308 100644
--- a/src/writer/hlsl/generator_impl_binary_test.cc
+++ b/src/writer/hlsl/generator_impl_binary_test.cc
@@ -60,8 +60,8 @@
 TEST_P(HlslBinaryTest, Emit_f32) {
   auto params = GetParam();
 
-  Global("left", ast::StorageClass::kFunction, ty.f32());
-  Global("right", ast::StorageClass::kFunction, ty.f32());
+  Global("left", ty.f32(), ast::StorageClass::kFunction);
+  Global("right", ty.f32(), ast::StorageClass::kFunction);
 
   auto* left = Expr("left");
   auto* right = Expr("right");
@@ -78,8 +78,8 @@
 TEST_P(HlslBinaryTest, Emit_u32) {
   auto params = GetParam();
 
-  Global("left", ast::StorageClass::kFunction, ty.u32());
-  Global("right", ast::StorageClass::kFunction, ty.u32());
+  Global("left", ty.u32(), ast::StorageClass::kFunction);
+  Global("right", ty.u32(), ast::StorageClass::kFunction);
 
   auto* left = Expr("left");
   auto* right = Expr("right");
@@ -96,8 +96,8 @@
 TEST_P(HlslBinaryTest, Emit_i32) {
   auto params = GetParam();
 
-  Global("left", ast::StorageClass::kFunction, ty.i32());
-  Global("right", ast::StorageClass::kFunction, ty.i32());
+  Global("left", ty.i32(), ast::StorageClass::kFunction);
+  Global("right", ty.i32(), ast::StorageClass::kFunction);
 
   auto* left = Expr("left");
   auto* right = Expr("right");
@@ -167,7 +167,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_MatrixScalar) {
-  Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = Expr("mat");
   auto* rhs = Expr(1.f);
 
@@ -182,7 +182,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_ScalarMatrix) {
-  Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = Expr(1.f);
   auto* rhs = Expr("mat");
 
@@ -197,7 +197,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_MatrixVector) {
-  Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = Expr("mat");
   auto* rhs = vec3<f32>(1.f, 1.f, 1.f);
 
@@ -212,7 +212,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_VectorMatrix) {
-  Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
   auto* rhs = Expr("mat");
 
@@ -227,7 +227,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Binary, Multiply_MatrixMatrix) {
-  Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = Expr("mat");
   auto* rhs = Expr("mat");
 
@@ -427,7 +427,7 @@
   auto* d = Expr("d");
 
   auto* var = Var(
-      "a", ast::StorageClass::kFunction, ty.bool_(),
+      "a", ty.bool_(), ast::StorageClass::kFunction,
       create<ast::BinaryExpression>(
           ast::BinaryOp::kLogicalOr,
           create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd, b, c), d),
@@ -483,10 +483,10 @@
 
   Func("foo", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
-  Global("a", ast::StorageClass::kNone, ty.bool_());
-  Global("b", ast::StorageClass::kNone, ty.bool_());
-  Global("c", ast::StorageClass::kNone, ty.bool_());
-  Global("d", ast::StorageClass::kNone, ty.bool_());
+  Global("a", ty.bool_(), ast::StorageClass::kNone);
+  Global("b", ty.bool_(), ast::StorageClass::kNone);
+  Global("c", ty.bool_(), ast::StorageClass::kNone);
+  Global("d", ty.bool_(), ast::StorageClass::kNone);
 
   ast::ExpressionList params;
   params.push_back(create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
diff --git a/src/writer/hlsl/generator_impl_call_test.cc b/src/writer/hlsl/generator_impl_call_test.cc
index 8094390..2ca0294 100644
--- a/src/writer/hlsl/generator_impl_call_test.cc
+++ b/src/writer/hlsl/generator_impl_call_test.cc
@@ -45,8 +45,8 @@
 TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
-  Global("param1", ast::StorageClass::kNone, ty.f32());
-  Global("param2", ast::StorageClass::kNone, ty.f32());
+  Global("param1", ty.f32(), ast::StorageClass::kNone);
+  Global("param2", ty.f32(), ast::StorageClass::kNone);
 
   auto* call = Call("my_func", "param1", "param2");
   WrapInFunction(call);
@@ -60,8 +60,8 @@
 TEST_F(HlslGeneratorImplTest_Call, EmitStatement_Call) {
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
-  Global("param1", ast::StorageClass::kNone, ty.f32());
-  Global("param2", ast::StorageClass::kNone, ty.f32());
+  Global("param1", ty.f32(), ast::StorageClass::kNone);
+  Global("param2", ty.f32(), ast::StorageClass::kNone);
 
   auto* call = create<ast::CallStatement>(Call("my_func", "param1", "param2"));
   WrapInFunction(call);
diff --git a/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc b/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
index fb07fa6..b2b6c94 100644
--- a/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
+++ b/src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
@@ -44,12 +44,12 @@
   //   int bar : TEXCOORD1;
   // };
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -87,12 +87,12 @@
   //   int bar : TEXCOORD1;
   // };
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -130,12 +130,12 @@
   //   int bar : TEXCOORD1;
   // };
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -173,12 +173,12 @@
   //   int bar : SV_Target1;
   // };
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -213,12 +213,12 @@
   //
   // -> Error, not allowed
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -248,12 +248,12 @@
   //
   // -> Error not allowed
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -289,12 +289,12 @@
   //   float depth : SV_Depth;
   // };
 
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord),
          });
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
          });
diff --git a/src/writer/hlsl/generator_impl_function_test.cc b/src/writer/hlsl/generator_impl_function_test.cc
index 3871aee..dceacf9 100644
--- a/src/writer/hlsl/generator_impl_function_test.cc
+++ b/src/writer/hlsl/generator_impl_function_test.cc
@@ -92,8 +92,8 @@
 
 TEST_F(HlslGeneratorImplTest_Function, Emit_Function_WithParams) {
   Func("my_func",
-       ast::VariableList{Var("a", ast::StorageClass::kNone, ty.f32()),
-                         Var("b", ast::StorageClass::kNone, ty.i32())},
+       ast::VariableList{Var("a", ty.f32(), ast::StorageClass::kNone),
+                         Var("b", ty.i32(), ast::StorageClass::kNone)},
        ty.void_(),
        ast::StatementList{
            create<ast::ReturnStatement>(),
@@ -132,12 +132,12 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_EntryPoint_NoReturn_InOut) {
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -172,12 +172,12 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_EntryPoint_WithInOutVars) {
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -213,12 +213,12 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_EntryPoint_WithInOut_Builtins) {
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord),
          });
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
          });
@@ -255,13 +255,13 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_EntryPoint_With_Uniform) {
-  Global("coord", ast::StorageClass::kUniform, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kUniform, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
          });
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "x"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -296,7 +296,7 @@
 
   auto* s = ty.struct_("Uniforms", str);
 
-  Global("uniforms", ast::StorageClass::kUniform, s, nullptr,
+  Global("uniforms", s, ast::StorageClass::kUniform, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
@@ -304,7 +304,7 @@
 
   AST().AddConstructedType(s);
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   create<ast::MemberAccessorExpression>(
                       MemberAccessor("uniforms", "coord"), Expr("x")),
                   ast::VariableDecorationList{});
@@ -345,13 +345,13 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
          });
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "b"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -386,13 +386,13 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kReadOnly, s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
          });
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "b"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -427,7 +427,7 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kWriteOnly, s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
@@ -466,7 +466,7 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
@@ -498,23 +498,23 @@
 TEST_F(
     HlslGeneratorImplTest_Function,
     Emit_FunctionDecoration_Called_By_EntryPoints_WithLocationGlobals_And_Params) {  // NOLINT
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
 
-  Global("val", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("val", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(0),
          });
 
   Func("sub_func",
-       ast::VariableList{Var("param", ast::StorageClass::kFunction, ty.f32())},
+       ast::VariableList{Var("param", ty.f32(), ast::StorageClass::kFunction)},
        ty.f32(),
        ast::StatementList{
            create<ast::AssignmentStatement>(Expr("bar"), Expr("foo")),
@@ -562,13 +562,13 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_Called_By_EntryPoints_NoUsedGlobals) {
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
          });
 
   Func("sub_func",
-       ast::VariableList{Var("param", ast::StorageClass::kFunction, ty.f32())},
+       ast::VariableList{Var("param", ty.f32(), ast::StorageClass::kFunction)},
        ty.f32(),
        ast::StatementList{
            create<ast::ReturnStatement>(Expr("param")),
@@ -608,18 +608,18 @@
 TEST_F(
     HlslGeneratorImplTest_Function,
     Emit_FunctionDecoration_Called_By_EntryPoints_WithBuiltinGlobals_And_Params) {  // NOLINT
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord),
          });
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
          });
 
   Func("sub_func",
-       ast::VariableList{Var("param", ast::StorageClass::kFunction, ty.f32())},
+       ast::VariableList{Var("param", ty.f32(), ast::StorageClass::kFunction)},
        ty.f32(),
        ast::StatementList{
            create<ast::AssignmentStatement>(Expr("depth"),
@@ -665,21 +665,21 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_Called_By_EntryPoint_With_Uniform) {
-  Global("coord", ast::StorageClass::kUniform, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kUniform, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
          });
 
   Func("sub_func",
-       ast::VariableList{Var("param", ast::StorageClass::kFunction, ty.f32())},
+       ast::VariableList{Var("param", ty.f32(), ast::StorageClass::kFunction)},
        ty.f32(),
        ast::StatementList{
            create<ast::ReturnStatement>(MemberAccessor("coord", "x")),
        },
        ast::FunctionDecorationList{});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   Call("sub_func", 1.0f), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -713,21 +713,21 @@
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_Called_By_EntryPoint_With_StorageBuffer) {
   type::AccessControl ac(ast::AccessControl::kReadWrite, ty.vec4<f32>());
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(1),
          });
 
   Func("sub_func",
-       ast::VariableList{Var("param", ast::StorageClass::kFunction, ty.f32())},
+       ast::VariableList{Var("param", ty.f32(), ast::StorageClass::kFunction)},
        ty.f32(),
        ast::StatementList{
            create<ast::ReturnStatement>(MemberAccessor("coord", "x")),
        },
        ast::FunctionDecorationList{});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   Call("sub_func", 1.0f), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -758,7 +758,7 @@
 
 TEST_F(HlslGeneratorImplTest_Function,
        Emit_FunctionDecoration_EntryPoints_WithGlobal_Nested_Return) {
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::LocationDecoration>(1),
          });
@@ -862,7 +862,7 @@
 TEST_F(HlslGeneratorImplTest_Function, Emit_Function_WithArrayParams) {
   Func(
       "my_func",
-      ast::VariableList{Var("a", ast::StorageClass::kNone, ty.array<f32, 5>())},
+      ast::VariableList{Var("a", ty.array<f32, 5>(), ast::StorageClass::kNone)},
       ty.void_(),
       ast::StatementList{
           create<ast::ReturnStatement>(),
@@ -908,14 +908,14 @@
 
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
 
-  Global("data", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("data", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(0),
          });
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("a", ast::VariableList{}, ty.void_(),
@@ -929,7 +929,7 @@
   }
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("b", ast::VariableList{}, ty.void_(),
diff --git a/src/writer/hlsl/generator_impl_import_test.cc b/src/writer/hlsl/generator_impl_import_test.cc
index c319f41..ad27e40 100644
--- a/src/writer/hlsl/generator_impl_import_test.cc
+++ b/src/writer/hlsl/generator_impl_import_test.cc
@@ -242,7 +242,7 @@
                          testing::Values(HlslImportData{"clamp", "clamp"}));
 
 TEST_F(HlslGeneratorImplTest_Import, HlslImportData_Determinant) {
-  Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("var", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = Call("determinant", "var");
   WrapInFunction(expr);
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc
index 4ec2683..d1353b9 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_test.cc
+++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -165,11 +165,11 @@
   ASSERT_NE(nullptr, call) << "Unhandled intrinsic";
   WrapInFunction(call);
 
-  Global("f2", ast::StorageClass::kFunction, ty.vec2<float>());
-  Global("f3", ast::StorageClass::kFunction, ty.vec3<float>());
-  Global("u2", ast::StorageClass::kFunction, ty.vec2<unsigned int>());
-  Global("b2", ast::StorageClass::kFunction, ty.vec2<bool>());
-  Global("m2x2", ast::StorageClass::kFunction, ty.mat2x2<float>());
+  Global("f2", ty.vec2<float>(), ast::StorageClass::kFunction);
+  Global("f3", ty.vec3<float>(), ast::StorageClass::kFunction);
+  Global("u2", ty.vec2<unsigned int>(), ast::StorageClass::kFunction);
+  Global("b2", ty.vec2<bool>(), ast::StorageClass::kFunction);
+  Global("m2x2", ty.mat2x2<float>(), ast::StorageClass::kFunction);
 
   GeneratorImpl& gen = Build();
 
@@ -264,8 +264,8 @@
 TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) {
   auto* call = Call("dot", "param1", "param2");
 
-  Global("param1", ast::StorageClass::kFunction, ty.vec3<f32>());
-  Global("param2", ast::StorageClass::kFunction, ty.vec3<f32>());
+  Global("param1", ty.vec3<f32>(), ast::StorageClass::kFunction);
+  Global("param2", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   WrapInFunction(call);
 
@@ -278,7 +278,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Snorm) {
   auto* call = Call("pack4x8snorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("p1", ty.vec4<f32>(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -292,7 +292,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Unorm) {
   auto* call = Call("pack4x8unorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("p1", ty.vec4<f32>(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -306,7 +306,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Snorm) {
   auto* call = Call("pack2x16snorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.vec2<f32>());
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -319,7 +319,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Unorm) {
   auto* call = Call("pack2x16unorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.vec2<f32>());
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -332,7 +332,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Pack2x16Float) {
   auto* call = Call("pack2x16float", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.vec2<f32>());
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -344,7 +344,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack4x8Snorm) {
   auto* call = Call("unpack4x8snorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -360,7 +360,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack4x8Unorm) {
   auto* call = Call("unpack4x8unorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -376,7 +376,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Snorm) {
   auto* call = Call("unpack2x16snorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -393,7 +393,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Unorm) {
   auto* call = Call("unpack2x16unorm", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
@@ -409,7 +409,7 @@
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Float) {
   auto* call = Call("unpack2x16float", "p1");
-  Global("p1", ast::StorageClass::kPrivate, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kPrivate);
   WrapInFunction(call);
   GeneratorImpl& gen = Build();
 
diff --git a/src/writer/hlsl/generator_impl_loop_test.cc b/src/writer/hlsl/generator_impl_loop_test.cc
index 480aeef..c4ab5c4 100644
--- a/src/writer/hlsl/generator_impl_loop_test.cc
+++ b/src/writer/hlsl/generator_impl_loop_test.cc
@@ -84,8 +84,8 @@
 }
 
 TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopNestedWithContinuing) {
-  Global("lhs", ast::StorageClass::kNone, ty.f32());
-  Global("rhs", ast::StorageClass::kNone, ty.f32());
+  Global("lhs", ty.f32(), ast::StorageClass::kNone);
+  Global("rhs", ty.f32(), ast::StorageClass::kNone);
 
   auto* body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::DiscardStatement>(),
@@ -160,15 +160,15 @@
   //   }
   // }
 
-  Global("rhs", ast::StorageClass::kNone, ty.f32());
+  Global("rhs", ty.f32(), ast::StorageClass::kNone);
 
-  auto* var = Var("lhs", ast::StorageClass::kFunction, ty.f32(), Expr(2.4f),
+  auto* var = Var("lhs", ty.f32(), ast::StorageClass::kFunction, Expr(2.4f),
                   ast::VariableDecorationList{});
 
   auto* body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(var),
       create<ast::VariableDeclStatement>(
-          Var("other", ast::StorageClass::kFunction, ty.f32())),
+          Var("other", ty.f32(), ast::StorageClass::kFunction)),
   });
 
   auto* lhs = Expr("lhs");
diff --git a/src/writer/hlsl/generator_impl_member_accessor_test.cc b/src/writer/hlsl/generator_impl_member_accessor_test.cc
index 37ad4fc..c1ee1c5 100644
--- a/src/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -45,7 +45,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Str", strct);
-  auto* str_var = Global("str", ast::StorageClass::kPrivate, s);
+  auto* str_var = Global("str", s, ast::StorageClass::kPrivate);
 
   auto* expr = MemberAccessor("str", "mem");
   WrapInFunction(expr);
@@ -75,7 +75,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "b");
   WrapInFunction(expr);
@@ -104,7 +104,7 @@
                             Member("b", ty.f32(), {MemberOffset(4)})},
       ast::StructDecorationList{});
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "a");
   WrapInFunction(expr);
@@ -137,8 +137,8 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* b_var = Global("b", ast::StorageClass::kPrivate, ty.mat2x3<f32>());
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* b_var = Global("b", ty.mat2x3<f32>(), ast::StorageClass::kPrivate);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor("data", "a");
   auto* rhs = Expr("b");
@@ -178,7 +178,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor("data", "a");
   auto* rhs = Construct(ty.mat2x3<f32>(), ast::ExpressionList{});
@@ -217,7 +217,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "a");
   WrapInFunction(expr);
@@ -255,7 +255,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "a");
   WrapInFunction(expr);
@@ -286,7 +286,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "a");
   WrapInFunction(expr);
@@ -318,7 +318,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = IndexAccessor(
       IndexAccessor(MemberAccessor("data", "a"), Expr(2)), Expr(1));
@@ -350,7 +350,7 @@
       ast::StructMemberList{Member("a", &ary, {MemberOffset(0)})},
       ast::StructDecorationList{});
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = IndexAccessor(MemberAccessor("data", "a"), Expr(2));
   WrapInFunction(expr);
@@ -381,7 +381,7 @@
       ast::StructMemberList{Member("a", &ary, {MemberOffset(0)})},
       ast::StructDecorationList{});
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = IndexAccessor(MemberAccessor("data", "a"),
                              Sub(Add(Expr(2), Expr(4)), Expr(3)));
@@ -412,7 +412,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor("data", "b");
   auto* rhs = Expr(2.0f);
@@ -448,7 +448,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = IndexAccessor(MemberAccessor("data", "a"), Expr(2));
   auto* rhs = Expr(2);
@@ -481,7 +481,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor("data", "a");
   auto* rhs = Expr(2);
@@ -514,7 +514,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor("data", "b");
   WrapInFunction(expr);
@@ -544,7 +544,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("Data", str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, s);
+  auto* coord_var = Global("data", s, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor("data", "b");
   auto* rhs = vec3<f32>(1.f, 2.f, 3.f);
@@ -596,7 +596,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* expr =
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b");
@@ -641,7 +641,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor(
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b"),
@@ -690,7 +690,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* expr = MemberAccessor(
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b"),
@@ -738,7 +738,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* expr = IndexAccessor(
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b"),
@@ -786,7 +786,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* lhs =
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b");
@@ -839,7 +839,7 @@
       ast::StructDecorationList{});
 
   auto* pre_struct = ty.struct_("Pre", pre_str);
-  auto* coord_var = Global("data", ast::StorageClass::kStorage, pre_struct);
+  auto* coord_var = Global("data", pre_struct, ast::StorageClass::kStorage);
 
   auto* lhs = MemberAccessor(
       MemberAccessor(IndexAccessor(MemberAccessor("data", "c"), Expr(2)), "b"),
@@ -862,7 +862,7 @@
 
 TEST_F(HlslGeneratorImplTest_MemberAccessor,
        EmitExpression_MemberAccessor_Swizzle_xyz) {
-  Global("my_vec", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = MemberAccessor("my_vec", "xyz");
   WrapInFunction(expr);
@@ -874,7 +874,7 @@
 
 TEST_F(HlslGeneratorImplTest_MemberAccessor,
        EmitExpression_MemberAccessor_Swizzle_gbr) {
-  Global("my_vec", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = MemberAccessor("my_vec", "gbr");
   WrapInFunction(expr);
diff --git a/src/writer/hlsl/generator_impl_module_constant_test.cc b/src/writer/hlsl/generator_impl_module_constant_test.cc
index d741433..0cc4430 100644
--- a/src/writer/hlsl/generator_impl_module_constant_test.cc
+++ b/src/writer/hlsl/generator_impl_module_constant_test.cc
@@ -34,7 +34,7 @@
 
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_ModuleConstant) {
   auto* var =
-      Const("pos", ast::StorageClass::kNone, ty.array<f32, 3>(),
+      Const("pos", ty.array<f32, 3>(), ast::StorageClass::kNone,
             array<f32, 3>(1.f, 2.f, 3.f), ast::VariableDecorationList{});
 
   GeneratorImpl& gen = Build();
@@ -44,7 +44,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant) {
-  auto* var = Const("pos", ast::StorageClass::kNone, ty.f32(), Expr(3.0f),
+  auto* var = Const("pos", ty.f32(), ast::StorageClass::kNone, Expr(3.0f),
                     ast::VariableDecorationList{
                         create<ast::ConstantIdDecoration>(23),
                     });
@@ -61,7 +61,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_ModuleConstant, Emit_SpecConstant_NoConstructor) {
-  auto* var = Const("pos", ast::StorageClass::kNone, ty.f32(), nullptr,
+  auto* var = Const("pos", ty.f32(), ast::StorageClass::kNone, nullptr,
                     ast::VariableDecorationList{
                         create<ast::ConstantIdDecoration>(23),
                     });
diff --git a/src/writer/hlsl/generator_impl_sanitizer_tests.cc b/src/writer/hlsl/generator_impl_sanitizer_tests.cc
index b0952ef..be4e677 100644
--- a/src/writer/hlsl/generator_impl_sanitizer_tests.cc
+++ b/src/writer/hlsl/generator_impl_sanitizer_tests.cc
@@ -32,7 +32,7 @@
 TEST_F(HlslSanitizerTest, PromoteArrayInitializerToConstVar) {
   auto* array_init = array<i32, 4>(1, 2, 3, 4);
   auto* array_index = IndexAccessor(array_init, 3);
-  auto* pos = Var("pos", ast::StorageClass::kFunction, ty.i32(), array_index);
+  auto* pos = Var("pos", ty.i32(), ast::StorageClass::kFunction, array_index);
 
   Func("main", ast::VariableList{}, ty.void_(),
        ast::StatementList{
diff --git a/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc b/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
index 8e14553..d784fa9 100644
--- a/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
+++ b/src/writer/hlsl/generator_impl_variable_decl_statement_test.cc
@@ -33,7 +33,7 @@
 using HlslGeneratorImplTest_VariableDecl = TestHelper;
 
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement) {
-  auto* var = Global("a", ast::StorageClass::kNone, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kNone);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -46,7 +46,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Const) {
-  auto* var = Const("a", ast::StorageClass::kNone, ty.f32());
+  auto* var = Const("a", ty.f32(), ast::StorageClass::kNone);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -59,7 +59,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Array) {
-  auto* var = Global("a", ast::StorageClass::kNone, ty.array<f32, 5>());
+  auto* var = Global("a", ty.array<f32, 5>(), ast::StorageClass::kNone);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -73,7 +73,7 @@
 
 TEST_F(HlslGeneratorImplTest_VariableDecl,
        Emit_VariableDeclStatement_Function) {
-  auto* var = Global("a", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kFunction);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -86,7 +86,7 @@
 }
 
 TEST_F(HlslGeneratorImplTest_VariableDecl, Emit_VariableDeclStatement_Private) {
-  auto* var = Global("a", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -100,8 +100,8 @@
 
 TEST_F(HlslGeneratorImplTest_VariableDecl,
        Emit_VariableDeclStatement_Initializer_Private) {
-  Global("initializer", ast::StorageClass::kNone, ty.f32());
-  auto* var = Global("a", ast::StorageClass::kPrivate, ty.f32(),
+  Global("initializer", ty.f32(), ast::StorageClass::kNone);
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kPrivate,
                      Expr("initializer"), ast::VariableDecorationList{});
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
@@ -115,7 +115,7 @@
 
 TEST_F(HlslGeneratorImplTest_VariableDecl,
        Emit_VariableDeclStatement_Initializer_ZeroVec) {
-  auto* var = Var("a", ast::StorageClass::kFunction, ty.vec3<f32>(),
+  auto* var = Var("a", ty.vec3<f32>(), ast::StorageClass::kFunction,
                   vec3<f32>(), ast::VariableDecorationList{});
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
@@ -130,7 +130,7 @@
 
 TEST_F(HlslGeneratorImplTest_VariableDecl,
        Emit_VariableDeclStatement_Initializer_ZeroMat) {
-  auto* var = Var("a", ast::StorageClass::kFunction, ty.mat2x3<f32>(),
+  auto* var = Var("a", ty.mat2x3<f32>(), ast::StorageClass::kFunction,
                   mat2x3<f32>(), ast::VariableDecorationList{});
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
diff --git a/src/writer/msl/generator_impl_call_test.cc b/src/writer/msl/generator_impl_call_test.cc
index 172f402..8a70461 100644
--- a/src/writer/msl/generator_impl_call_test.cc
+++ b/src/writer/msl/generator_impl_call_test.cc
@@ -47,8 +47,8 @@
 TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithParams) {
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
-  Global("param1", ast::StorageClass::kNone, ty.f32());
-  Global("param2", ast::StorageClass::kNone, ty.f32());
+  Global("param1", ty.f32(), ast::StorageClass::kNone);
+  Global("param2", ty.f32(), ast::StorageClass::kNone);
 
   auto* call = Call("my_func", "param1", "param2");
   WrapInFunction(call);
@@ -62,8 +62,8 @@
 TEST_F(MslGeneratorImplTest, EmitStatement_Call) {
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
-  Global("param1", ast::StorageClass::kNone, ty.f32());
-  Global("param2", ast::StorageClass::kNone, ty.f32());
+  Global("param1", ty.f32(), ast::StorageClass::kNone);
+  Global("param2", ty.f32(), ast::StorageClass::kNone);
 
   auto* call = Call("my_func", "param1", "param2");
   auto* stmt = create<ast::CallStatement>(call);
diff --git a/src/writer/msl/generator_impl_function_entry_point_data_test.cc b/src/writer/msl/generator_impl_function_entry_point_data_test.cc
index 301ad31..ada6284 100644
--- a/src/writer/msl/generator_impl_function_entry_point_data_test.cc
+++ b/src/writer/msl/generator_impl_function_entry_point_data_test.cc
@@ -47,10 +47,10 @@
   //   int bar [[attribute(1)]];
   // };
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -84,10 +84,10 @@
   //   int bar [[user(locn1)]];
   // };
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -121,10 +121,10 @@
   //   int bar [[user(locn1)]];
   // };
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -158,10 +158,10 @@
   //   int bar [[color(1)]];
   // };
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -192,10 +192,10 @@
   //
   // -> Error, not allowed
 
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kInput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -221,10 +221,10 @@
   //
   // -> Error not allowed
 
-  Global("foo", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.i32(), nullptr,
+  Global("bar", ty.i32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -255,11 +255,11 @@
   //   float depth [[depth(any)]];
   // };
 
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord)});
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth)});
 
diff --git a/src/writer/msl/generator_impl_function_test.cc b/src/writer/msl/generator_impl_function_test.cc
index 3a5d9b6..05d51cc 100644
--- a/src/writer/msl/generator_impl_function_test.cc
+++ b/src/writer/msl/generator_impl_function_test.cc
@@ -99,8 +99,8 @@
 
 TEST_F(MslGeneratorImplTest, Emit_Function_WithParams) {
   ast::VariableList params;
-  params.push_back(Var("a", ast::StorageClass::kNone, ty.f32()));
-  params.push_back(Var("b", ast::StorageClass::kNone, ty.i32()));
+  params.push_back(Var("a", ty.f32(), ast::StorageClass::kNone));
+  params.push_back(Var("b", ty.i32(), ast::StorageClass::kNone));
 
   Func("my_func", params, ty.void_(),
        ast::StatementList{
@@ -142,10 +142,10 @@
 
 TEST_F(MslGeneratorImplTest,
        Emit_FunctionDecoration_EntryPoint_NoReturn_InOut) {
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   Func("main", ast::VariableList{}, ty.void_(),
@@ -178,10 +178,10 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_FunctionDecoration_EntryPoint_WithInOutVars) {
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto body = ast::StatementList{
@@ -216,11 +216,11 @@
 
 TEST_F(MslGeneratorImplTest,
        Emit_FunctionDecoration_EntryPoint_WithInOut_Builtins) {
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord)});
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth)});
 
@@ -254,11 +254,11 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_FunctionDecoration_EntryPoint_With_Uniform) {
-  Global("coord", ast::StorageClass::kUniform, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kUniform, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "x"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -295,11 +295,11 @@
 
   AST().AddConstructedType(s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "b"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -340,11 +340,11 @@
   type::AccessControl ac(ast::AccessControl::kReadOnly, s);
   AST().AddConstructedType(s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   MemberAccessor("coord", "b"), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -377,17 +377,17 @@
 TEST_F(
     MslGeneratorImplTest,
     Emit_FunctionDecoration_Called_By_EntryPoints_WithLocationGlobals_And_Params) {  // NOLINT
-  Global("foo", ast::StorageClass::kInput, ty.f32(), nullptr,
+  Global("foo", ty.f32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
-  Global("val", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("val", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(0)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   auto body = ast::StatementList{
       create<ast::AssignmentStatement>(Expr("bar"), Expr("foo")),
@@ -437,12 +437,12 @@
 
 TEST_F(MslGeneratorImplTest,
        Emit_FunctionDecoration_Called_By_EntryPoints_NoUsedGlobals) {
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   Func("sub_func", params, ty.f32(),
        ast::StatementList{
@@ -485,16 +485,16 @@
 TEST_F(
     MslGeneratorImplTest,
     Emit_FunctionDecoration_Called_By_EntryPoints_WithBuiltinGlobals_And_Params) {  // NOLINT
-  Global("coord", ast::StorageClass::kInput, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragCoord)});
 
-  Global("depth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("depth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   auto body = ast::StatementList{
       create<ast::AssignmentStatement>(Expr("depth"),
@@ -539,12 +539,12 @@
 
 TEST_F(MslGeneratorImplTest,
        Emit_FunctionDecoration_Called_By_EntryPoint_With_Uniform) {
-  Global("coord", ast::StorageClass::kUniform, ty.vec4<f32>(), nullptr,
+  Global("coord", ty.vec4<f32>(), ast::StorageClass::kUniform, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   auto body = ast::StatementList{
       create<ast::ReturnStatement>(MemberAccessor("coord", "x")),
@@ -555,7 +555,7 @@
   ast::ExpressionList expr;
   expr.push_back(Expr(1.0f));
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   Call("sub_func", 1.0f), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -595,19 +595,19 @@
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
   AST().AddConstructedType(s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   auto body = ast::StatementList{
       create<ast::ReturnStatement>(MemberAccessor("coord", "b"))};
 
   Func("sub_func", params, ty.f32(), body, ast::FunctionDecorationList{});
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   Call("sub_func", 1.0f), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -652,12 +652,12 @@
   type::AccessControl ac(ast::AccessControl::kReadOnly, s);
   AST().AddConstructedType(s);
 
-  Global("coord", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("coord", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(1)});
 
   ast::VariableList params;
-  params.push_back(Var("param", ast::StorageClass::kFunction, ty.f32()));
+  params.push_back(Var("param", ty.f32(), ast::StorageClass::kFunction));
 
   auto body = ast::StatementList{
       create<ast::ReturnStatement>(MemberAccessor("coord", "b"))};
@@ -667,7 +667,7 @@
   ast::ExpressionList expr;
   expr.push_back(Expr(1.0f));
 
-  auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+  auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                   Call("sub_func", 1.0f), ast::VariableDecorationList{});
 
   Func("frag_main", ast::VariableList{}, ty.void_(),
@@ -703,7 +703,7 @@
 
 TEST_F(MslGeneratorImplTest,
        Emit_FunctionDecoration_EntryPoints_WithGlobal_Nested_Return) {
-  Global("bar", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("bar", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{create<ast::LocationDecoration>(1)});
 
   auto* list = create<ast::BlockStatement>(ast::StatementList{
@@ -765,7 +765,7 @@
 
 TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) {
   ast::VariableList params;
-  params.push_back(Var("a", ast::StorageClass::kNone, ty.array<f32, 5>()));
+  params.push_back(Var("a", ty.array<f32, 5>(), ast::StorageClass::kNone));
 
   Func("my_func", params, ty.void_(),
        ast::StatementList{
@@ -814,14 +814,14 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
 
-  Global("data", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("data", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{create<ast::BindingDecoration>(0),
                                      create<ast::GroupDecoration>(0)});
 
   AST().AddConstructedType(s);
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("a", ast::VariableList{}, ty.void_(),
@@ -835,7 +835,7 @@
   }
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("b", ast::VariableList{}, ty.void_(),
diff --git a/src/writer/msl/generator_impl_import_test.cc b/src/writer/msl/generator_impl_import_test.cc
index 7b365f4..94530aa 100644
--- a/src/writer/msl/generator_impl_import_test.cc
+++ b/src/writer/msl/generator_impl_import_test.cc
@@ -204,7 +204,7 @@
                                          MslImportData{"clamp", "clamp"}));
 
 TEST_F(MslGeneratorImplTest, MslImportData_Determinant) {
-  Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  Global("var", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = Call("determinant", "var");
 
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc
index 3873b4c..059c9b5 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -176,13 +176,13 @@
   ASSERT_NE(nullptr, call) << "Unhandled intrinsic";
   WrapInFunction(call);
 
-  Global("f2", ast::StorageClass::kFunction, ty.vec2<float>());
-  Global("f3", ast::StorageClass::kFunction, ty.vec3<float>());
-  Global("f4", ast::StorageClass::kFunction, ty.vec4<float>());
-  Global("u1", ast::StorageClass::kFunction, ty.u32());
-  Global("u2", ast::StorageClass::kFunction, ty.vec2<unsigned int>());
-  Global("b2", ast::StorageClass::kFunction, ty.vec2<bool>());
-  Global("m2x2", ast::StorageClass::kFunction, ty.mat2x2<float>());
+  Global("f2", ty.vec2<float>(), ast::StorageClass::kFunction);
+  Global("f3", ty.vec3<float>(), ast::StorageClass::kFunction);
+  Global("f4", ty.vec4<float>(), ast::StorageClass::kFunction);
+  Global("u1", ty.u32(), ast::StorageClass::kFunction);
+  Global("u2", ty.vec2<unsigned int>(), ast::StorageClass::kFunction);
+  Global("b2", ty.vec2<bool>(), ast::StorageClass::kFunction);
+  Global("m2x2", ty.mat2x2<float>(), ast::StorageClass::kFunction);
 
   GeneratorImpl& gen = Build();
 
@@ -293,8 +293,8 @@
                       "metal::unpack_unorm2x16_to_float"}));
 
 TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
-  Global("param1", ast::StorageClass::kFunction, ty.vec2<f32>());
-  Global("param2", ast::StorageClass::kFunction, ty.vec2<f32>());
+  Global("param1", ty.vec2<f32>(), ast::StorageClass::kFunction);
+  Global("param2", ty.vec2<f32>(), ast::StorageClass::kFunction);
 
   auto* call = Call("dot", "param1", "param2");
   WrapInFunction(call);
@@ -308,7 +308,7 @@
 
 TEST_F(MslGeneratorImplTest, Pack2x16Float) {
   auto* call = Call("pack2x16float", "p1");
-  Global("p1", ast::StorageClass::kFunction, ty.vec2<f32>());
+  Global("p1", ty.vec2<f32>(), ast::StorageClass::kFunction);
   WrapInFunction(call);
 
   GeneratorImpl& gen = Build();
@@ -320,7 +320,7 @@
 
 TEST_F(MslGeneratorImplTest, Unpack2x16Float) {
   auto* call = Call("unpack2x16float", "p1");
-  Global("p1", ast::StorageClass::kFunction, ty.u32());
+  Global("p1", ty.u32(), ast::StorageClass::kFunction);
   WrapInFunction(call);
 
   GeneratorImpl& gen = Build();
diff --git a/src/writer/msl/generator_impl_loop_test.cc b/src/writer/msl/generator_impl_loop_test.cc
index 9741d4f..53da280 100644
--- a/src/writer/msl/generator_impl_loop_test.cc
+++ b/src/writer/msl/generator_impl_loop_test.cc
@@ -84,8 +84,8 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_LoopNestedWithContinuing) {
-  Global("lhs", ast::StorageClass::kNone, ty.f32());
-  Global("rhs", ast::StorageClass::kNone, ty.f32());
+  Global("lhs", ty.f32(), ast::StorageClass::kNone);
+  Global("rhs", ty.f32(), ast::StorageClass::kNone);
 
   auto* body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::DiscardStatement>(),
@@ -157,15 +157,15 @@
   //   }
   // }
 
-  Global("rhs", ast::StorageClass::kNone, ty.f32());
+  Global("rhs", ty.f32(), ast::StorageClass::kNone);
 
-  auto* var = Var("lhs", ast::StorageClass::kFunction, ty.f32(), Expr(2.4f),
+  auto* var = Var("lhs", ty.f32(), ast::StorageClass::kFunction, Expr(2.4f),
                   ast::VariableDecorationList{});
 
   auto* body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(var),
       create<ast::VariableDeclStatement>(
-          Var("other", ast::StorageClass::kFunction, ty.f32()))});
+          Var("other", ty.f32(), ast::StorageClass::kFunction))});
 
   auto* continuing = create<ast::BlockStatement>(ast::StatementList{
       create<ast::AssignmentStatement>(Expr("lhs"), Expr("rhs")),
diff --git a/src/writer/msl/generator_impl_member_accessor_test.cc b/src/writer/msl/generator_impl_member_accessor_test.cc
index 6370bc4..d712c75 100644
--- a/src/writer/msl/generator_impl_member_accessor_test.cc
+++ b/src/writer/msl/generator_impl_member_accessor_test.cc
@@ -29,12 +29,13 @@
 using MslGeneratorImplTest = TestHelper;
 
 TEST_F(MslGeneratorImplTest, EmitExpression_MemberAccessor) {
-  Global("str", ast::StorageClass::kPrivate,
+  Global("str",
          ty.struct_("my_str", create<ast::Struct>(
                                   ast::StructMemberList{
                                       Member("mem", ty.f32()),
                                   },
-                                  ast::StructDecorationList{})));
+                                  ast::StructDecorationList{})),
+         ast::StorageClass::kPrivate);
   auto* expr = MemberAccessor("str", "mem");
   WrapInFunction(expr);
 
@@ -45,7 +46,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, EmitExpression_MemberAccessor_Swizzle_xyz) {
-  Global("my_vec", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = MemberAccessor("my_vec", "xyz");
   WrapInFunction(expr);
@@ -56,7 +57,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, EmitExpression_MemberAccessor_Swizzle_gbr) {
-  Global("my_vec", ast::StorageClass::kPrivate, ty.vec4<f32>());
+  Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = MemberAccessor("my_vec", "gbr");
   WrapInFunction(expr);
diff --git a/src/writer/msl/generator_impl_module_constant_test.cc b/src/writer/msl/generator_impl_module_constant_test.cc
index be91b0a..5ad0da7 100644
--- a/src/writer/msl/generator_impl_module_constant_test.cc
+++ b/src/writer/msl/generator_impl_module_constant_test.cc
@@ -36,7 +36,7 @@
 
 TEST_F(MslGeneratorImplTest, Emit_ModuleConstant) {
   auto* var =
-      Const("pos", ast::StorageClass::kNone, ty.array<f32, 3>(),
+      Const("pos", ty.array<f32, 3>(), ast::StorageClass::kNone,
             array<f32, 3>(1.f, 2.f, 3.f), ast::VariableDecorationList{});
 
   GeneratorImpl& gen = Build();
@@ -46,7 +46,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_SpecConstant) {
-  auto* var = Const("pos", ast::StorageClass::kNone, ty.f32(), Expr(3.f),
+  auto* var = Const("pos", ty.f32(), ast::StorageClass::kNone, Expr(3.f),
                     ast::VariableDecorationList{
                         create<ast::ConstantIdDecoration>(23),
                     });
diff --git a/src/writer/msl/generator_impl_type_test.cc b/src/writer/msl/generator_impl_type_test.cc
index c223f0c..9321bba 100644
--- a/src/writer/msl/generator_impl_type_test.cc
+++ b/src/writer/msl/generator_impl_type_test.cc
@@ -429,7 +429,7 @@
       create<type::AccessControl>(params.ro ? ast::AccessControl::kReadOnly
                                             : ast::AccessControl::kWriteOnly,
                                   s);
-  Global("test_var", ast::StorageClass::kNone, ac);
+  Global("test_var", ac, ast::StorageClass::kNone);
 
   GeneratorImpl& gen = Build();
 
diff --git a/src/writer/msl/generator_impl_variable_decl_statement_test.cc b/src/writer/msl/generator_impl_variable_decl_statement_test.cc
index 33f230e..0931495 100644
--- a/src/writer/msl/generator_impl_variable_decl_statement_test.cc
+++ b/src/writer/msl/generator_impl_variable_decl_statement_test.cc
@@ -41,7 +41,7 @@
 using MslGeneratorImplTest = TestHelper;
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement) {
-  auto* var = Var("a", ast::StorageClass::kNone, ty.f32());
+  auto* var = Var("a", ty.f32(), ast::StorageClass::kNone);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
 
@@ -54,7 +54,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Const) {
-  auto* var = Const("a", ast::StorageClass::kNone, ty.f32());
+  auto* var = Const("a", ty.f32(), ast::StorageClass::kNone);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
 
@@ -69,7 +69,7 @@
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Array) {
   type::Array ary(ty.f32(), 5, ast::ArrayDecorationList{});
 
-  auto* var = Var("a", ast::StorageClass::kNone, &ary);
+  auto* var = Var("a", &ary, ast::StorageClass::kNone);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
 
@@ -88,7 +88,7 @@
       ast::StructDecorationList{});
 
   auto* s = ty.struct_("S", str);
-  auto* var = Var("a", ast::StorageClass::kNone, s);
+  auto* var = Var("a", s, ast::StorageClass::kNone);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
 
@@ -102,7 +102,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Vector) {
-  auto* var = Var("a", ast::StorageClass::kFunction, ty.vec2<f32>());
+  auto* var = Var("a", ty.vec2<f32>(), ast::StorageClass::kFunction);
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
 
@@ -115,7 +115,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Matrix) {
-  auto* var = Var("a", ast::StorageClass::kFunction, ty.mat3x2<f32>());
+  auto* var = Var("a", ty.mat3x2<f32>(), ast::StorageClass::kFunction);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
@@ -129,7 +129,7 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Private) {
-  auto* var = Global("a", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kPrivate);
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
   GeneratorImpl& gen = Build();
@@ -141,8 +141,8 @@
 }
 
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Initializer_Private) {
-  Global("initializer", ast::StorageClass::kNone, ty.f32());
-  auto* var = Global("a", ast::StorageClass::kPrivate, ty.f32(),
+  Global("initializer", ty.f32(), ast::StorageClass::kNone);
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kPrivate,
                      Expr("initializer"), ast::VariableDecorationList{});
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -156,7 +156,7 @@
 TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Initializer_ZeroVec) {
   auto* zero_vec = vec3<f32>();
 
-  auto* var = Var("a", ast::StorageClass::kFunction, ty.vec3<f32>(), zero_vec,
+  auto* var = Var("a", ty.vec3<f32>(), ast::StorageClass::kFunction, zero_vec,
                   ast::VariableDecorationList{});
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
diff --git a/src/writer/spirv/builder_accessor_expression_test.cc b/src/writer/spirv/builder_accessor_expression_test.cc
index c0c8ffb..40d5b5f 100644
--- a/src/writer/spirv/builder_accessor_expression_test.cc
+++ b/src/writer/spirv/builder_accessor_expression_test.cc
@@ -49,7 +49,7 @@
   // vec3<f32> ary;
   // ary[1]  -> ptr<f32>
 
-  auto* var = Global("ary", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ary", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* ary = Expr("ary");
   auto* idx_expr = Expr(1);
@@ -85,8 +85,8 @@
   // idx : i32;
   // ary[idx]  -> ptr<f32>
 
-  auto* var = Global("ary", ast::StorageClass::kFunction, ty.vec3<f32>());
-  auto* idx = Global("idx", ast::StorageClass::kFunction, ty.i32());
+  auto* var = Global("ary", ty.vec3<f32>(), ast::StorageClass::kFunction);
+  auto* idx = Global("idx", ty.i32(), ast::StorageClass::kFunction);
 
   auto* ary = Expr("ary");
   auto* idx_expr = Expr("idx");
@@ -125,7 +125,7 @@
   // vec3<f32> ary;
   // ary[1 + 2]  -> ptr<f32>
 
-  auto* var = Global("ary", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ary", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* ary = Expr("ary");
 
@@ -163,7 +163,7 @@
   // ary = array<vec3<f32>, 4>
   // ary[3][2];
 
-  auto* var = Global("ary", ast::StorageClass::kFunction, &ary4);
+  auto* var = Global("ary", &ary4, ast::StorageClass::kFunction);
 
   auto* expr = IndexAccessor(IndexAccessor("ary", 3), 2);
   WrapInFunction(expr);
@@ -201,7 +201,7 @@
   // var a : array<vec3<f32>, 4>;
   // a[2].xy;
 
-  auto* var = Global("ary", ast::StorageClass::kFunction, &ary4);
+  auto* var = Global("ary", &ary4, ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor(IndexAccessor("ary", 2), "xy");
   WrapInFunction(expr);
@@ -247,7 +247,7 @@
       ast::StructDecorationList{});
 
   auto* s_type = ty.struct_("my_struct", s);
-  auto* var = Global("ident", ast::StorageClass::kFunction, s_type);
+  auto* var = Global("ident", s_type, ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor("ident", "b");
   WrapInFunction(expr);
@@ -295,7 +295,7 @@
       create<ast::Struct>(ast::StructMemberList{Member("inner", inner_struct)},
                           ast::StructDecorationList{}));
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, s_type);
+  auto* var = Global("ident", s_type, ast::StorageClass::kFunction);
   auto* expr = MemberAccessor(MemberAccessor("ident", "inner"), "a");
   WrapInFunction(expr);
 
@@ -345,7 +345,7 @@
       create<ast::Struct>(ast::StructMemberList{Member("inner", alias)},
                           ast::StructDecorationList{}));
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, s_type);
+  auto* var = Global("ident", s_type, ast::StorageClass::kFunction);
   auto* expr = MemberAccessor(MemberAccessor("ident", "inner"), "a");
   WrapInFunction(expr);
 
@@ -393,7 +393,7 @@
       create<ast::Struct>(ast::StructMemberList{Member("inner", inner_struct)},
                           ast::StructDecorationList{}));
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, s_type);
+  auto* var = Global("ident", s_type, ast::StorageClass::kFunction);
   auto* expr = create<ast::AssignmentStatement>(
       MemberAccessor(MemberAccessor("ident", "inner"), "a"), Expr(2.0f));
   WrapInFunction(expr);
@@ -445,8 +445,8 @@
       create<ast::Struct>(ast::StructMemberList{Member("inner", inner_struct)},
                           ast::StructDecorationList{}));
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, s_type);
-  auto* store = Global("store", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Global("ident", s_type, ast::StorageClass::kFunction);
+  auto* store = Global("store", ty.f32(), ast::StorageClass::kFunction);
 
   auto* rhs = MemberAccessor(MemberAccessor("ident", "inner"), "a");
   auto* expr = create<ast::AssignmentStatement>(Expr("store"), rhs);
@@ -484,7 +484,7 @@
 TEST_F(BuilderTest, MemberAccessor_Swizzle_Single) {
   // ident.y
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ident", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor("ident", "y");
   WrapInFunction(expr);
@@ -515,7 +515,7 @@
 TEST_F(BuilderTest, MemberAccessor_Swizzle_MultipleNames) {
   // ident.yx
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ident", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor("ident", "yx");
   WrapInFunction(expr);
@@ -545,7 +545,7 @@
 TEST_F(BuilderTest, MemberAccessor_Swizzle_of_Swizzle) {
   // ident.yxz.xz
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ident", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor(MemberAccessor("ident", "yxz"), "xz");
   WrapInFunction(expr);
@@ -576,7 +576,7 @@
 TEST_F(BuilderTest, MemberAccessor_Member_of_Swizzle) {
   // ident.yxz.x
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ident", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = MemberAccessor(MemberAccessor("ident", "yxz"), "x");
   WrapInFunction(expr);
@@ -606,7 +606,7 @@
 TEST_F(BuilderTest, MemberAccessor_Array_of_Swizzle) {
   // index.yxz[1]
 
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("ident", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = IndexAccessor(MemberAccessor("ident", "yxz"), 1);
   WrapInFunction(expr);
@@ -662,7 +662,7 @@
   auto* a_type = ty.struct_("A", s);
 
   type::Array a_ary_type(a_type, 2, ast::ArrayDecorationList{});
-  auto* var = Global("index", ast::StorageClass::kFunction, &a_ary_type);
+  auto* var = Global("index", &a_ary_type, ast::StorageClass::kFunction);
   auto* expr = MemberAccessor(
       MemberAccessor(
           MemberAccessor(
@@ -719,7 +719,7 @@
   type::Array arr(ty.vec2<f32>(), 3, ast::ArrayDecorationList{});
 
   auto* var =
-      GlobalConst("pos", ast::StorageClass::kPrivate, &arr,
+      GlobalConst("pos", &arr, ast::StorageClass::kPrivate,
                   Construct(&arr, vec2<f32>(0.0f, 0.5f),
                             vec2<f32>(-0.5f, -0.5f), vec2<f32>(0.5f, -0.5f)),
                   ast::VariableDecorationList{});
@@ -763,7 +763,7 @@
   // const pos : vec2<f32> = vec2<f32>(0.0, 0.5);
   // pos[1]
 
-  auto* var = GlobalConst("pos", ast::StorageClass::kPrivate, ty.vec2<f32>(),
+  auto* var = GlobalConst("pos", ty.vec2<f32>(), ast::StorageClass::kPrivate,
                           vec2<f32>(0.0f, 0.5f), ast::VariableDecorationList{});
 
   auto* expr = IndexAccessor("pos", 1u);
diff --git a/src/writer/spirv/builder_assign_test.cc b/src/writer/spirv/builder_assign_test.cc
index 9cf0c01..282a9b5 100644
--- a/src/writer/spirv/builder_assign_test.cc
+++ b/src/writer/spirv/builder_assign_test.cc
@@ -42,7 +42,7 @@
 using BuilderTest = TestHelper;
 
 TEST_F(BuilderTest, Assign_Var) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput);
 
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), Expr(1.f));
 
@@ -70,7 +70,7 @@
 }
 
 TEST_F(BuilderTest, Assign_Var_OutsideFunction_IsError) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput);
 
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), Expr(1.f));
 
@@ -89,7 +89,7 @@
 }
 
 TEST_F(BuilderTest, Assign_Var_ZeroConstructor) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* val = vec3<f32>();
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), val);
@@ -120,7 +120,7 @@
 TEST_F(BuilderTest, Assign_Var_Complex_ConstructorWithExtract) {
   auto* init = vec3<f32>(vec2<f32>(1.f, 2.f), 3.f);
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), init);
 
@@ -157,7 +157,7 @@
 TEST_F(BuilderTest, Assign_Var_Complex_Constructor) {
   auto* init = vec3<f32>(1.f, 2.f, 3.f);
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), init);
 
@@ -200,7 +200,7 @@
       ast::StructDecorationList{});
 
   auto* s_type = ty.struct_("my_struct", s);
-  auto* v = Global("ident", ast::StorageClass::kFunction, s_type);
+  auto* v = Global("ident", s_type, ast::StorageClass::kFunction);
 
   auto* assign =
       create<ast::AssignmentStatement>(MemberAccessor("ident", "b"), Expr(4.f));
@@ -233,7 +233,7 @@
 }
 
 TEST_F(BuilderTest, Assign_Vector) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* val = vec3<f32>(1.f, 1.f, 3.f);
   auto* assign = create<ast::AssignmentStatement>(Expr("var"), val);
@@ -267,7 +267,7 @@
 TEST_F(BuilderTest, Assign_Vector_MemberByName) {
   // var.y = 1
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* assign =
       create<ast::AssignmentStatement>(MemberAccessor("var", "y"), Expr(1.f));
@@ -303,7 +303,7 @@
 TEST_F(BuilderTest, Assign_Vector_MemberByIndex) {
   // var[1] = 1
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.vec3<f32>());
+  auto* v = Global("var", ty.vec3<f32>(), ast::StorageClass::kOutput);
 
   auto* assign =
       create<ast::AssignmentStatement>(IndexAccessor("var", 1), Expr(1.f));
diff --git a/src/writer/spirv/builder_binary_expression_test.cc b/src/writer/spirv/builder_binary_expression_test.cc
index 97f0b54..7009de6 100644
--- a/src/writer/spirv/builder_binary_expression_test.cc
+++ b/src/writer/spirv/builder_binary_expression_test.cc
@@ -100,7 +100,7 @@
 TEST_P(BinaryArithSignedIntegerTest, Scalar_Loads) {
   auto param = GetParam();
 
-  auto* var = Global("param", ast::StorageClass::kFunction, ty.i32());
+  auto* var = Global("param", ty.i32(), ast::StorageClass::kFunction);
   auto* expr =
       create<ast::BinaryExpression>(param.op, Expr("param"), Expr("param"));
 
@@ -491,7 +491,7 @@
 }
 
 TEST_F(BuilderTest, Binary_Multiply_MatrixScalar) {
-  auto* var = Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  auto* var = Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kMultiply,
                                              Expr("mat"), Expr(1.f));
@@ -519,7 +519,7 @@
 }
 
 TEST_F(BuilderTest, Binary_Multiply_ScalarMatrix) {
-  auto* var = Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  auto* var = Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kMultiply,
                                              Expr(1.f), Expr("mat"));
@@ -547,7 +547,7 @@
 }
 
 TEST_F(BuilderTest, Binary_Multiply_MatrixVector) {
-  auto* var = Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  auto* var = Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* rhs = vec3<f32>(1.f, 1.f, 1.f);
 
   auto* expr =
@@ -577,7 +577,7 @@
 }
 
 TEST_F(BuilderTest, Binary_Multiply_VectorMatrix) {
-  auto* var = Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  auto* var = Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
   auto* lhs = vec3<f32>(1.f, 1.f, 1.f);
 
   auto* expr =
@@ -607,7 +607,7 @@
 }
 
 TEST_F(BuilderTest, Binary_Multiply_MatrixMatrix) {
-  auto* var = Global("mat", ast::StorageClass::kFunction, ty.mat3x3<f32>());
+  auto* var = Global("mat", ty.mat3x3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kMultiply,
                                              Expr("mat"), Expr("mat"));
@@ -674,9 +674,9 @@
 }
 
 TEST_F(BuilderTest, Binary_LogicalAnd_WithLoads) {
-  auto* a_var = Global("a", ast::StorageClass::kFunction, ty.bool_(),
+  auto* a_var = Global("a", ty.bool_(), ast::StorageClass::kFunction,
                        Expr(true), ast::VariableDecorationList{});
-  auto* b_var = Global("b", ast::StorageClass::kFunction, ty.bool_(),
+  auto* b_var = Global("b", ty.bool_(), ast::StorageClass::kFunction,
                        Expr(false), ast::VariableDecorationList{});
 
   auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
@@ -833,9 +833,9 @@
 }
 
 TEST_F(BuilderTest, Binary_LogicalOr_WithLoads) {
-  auto* a_var = Global("a", ast::StorageClass::kFunction, ty.bool_(),
+  auto* a_var = Global("a", ty.bool_(), ast::StorageClass::kFunction,
                        Expr(true), ast::VariableDecorationList{});
-  auto* b_var = Global("b", ast::StorageClass::kFunction, ty.bool_(),
+  auto* b_var = Global("b", ty.bool_(), ast::StorageClass::kFunction,
                        Expr(false), ast::VariableDecorationList{});
 
   auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kLogicalOr,
diff --git a/src/writer/spirv/builder_block_test.cc b/src/writer/spirv/builder_block_test.cc
index b19a750..ac6327b 100644
--- a/src/writer/spirv/builder_block_test.cc
+++ b/src/writer/spirv/builder_block_test.cc
@@ -39,11 +39,11 @@
   // serves to prove the block code is pushing new scopes as needed.
   auto* inner = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(
-          Var("var", ast::StorageClass::kFunction, ty.f32())),
+          Var("var", ty.f32(), ast::StorageClass::kFunction)),
       create<ast::AssignmentStatement>(Expr("var"), Expr(2.f))});
   auto* outer = create<ast::BlockStatement>(ast::StatementList{
       create<ast::VariableDeclStatement>(
-          Var("var", ast::StorageClass::kFunction, ty.f32())),
+          Var("var", ty.f32(), ast::StorageClass::kFunction)),
       create<ast::AssignmentStatement>(Expr("var"), Expr(1.f)), inner,
       create<ast::AssignmentStatement>(Expr("var"), Expr(3.f))});
 
diff --git a/src/writer/spirv/builder_call_test.cc b/src/writer/spirv/builder_call_test.cc
index 7f8db9c..21fc547 100644
--- a/src/writer/spirv/builder_call_test.cc
+++ b/src/writer/spirv/builder_call_test.cc
@@ -41,8 +41,8 @@
 
 TEST_F(BuilderTest, Expression_Call) {
   ast::VariableList func_params;
-  func_params.push_back(Var("a", ast::StorageClass::kFunction, ty.f32()));
-  func_params.push_back(Var("b", ast::StorageClass::kFunction, ty.f32()));
+  func_params.push_back(Var("a", ty.f32(), ast::StorageClass::kFunction));
+  func_params.push_back(Var("b", ty.f32(), ast::StorageClass::kFunction));
 
   auto* a_func =
       Func("a_func", func_params, ty.f32(),
@@ -90,8 +90,8 @@
 
 TEST_F(BuilderTest, Statement_Call) {
   ast::VariableList func_params;
-  func_params.push_back(Var("a", ast::StorageClass::kFunction, ty.f32()));
-  func_params.push_back(Var("b", ast::StorageClass::kFunction, ty.f32()));
+  func_params.push_back(Var("a", ty.f32(), ast::StorageClass::kFunction));
+  func_params.push_back(Var("b", ty.f32(), ast::StorageClass::kFunction));
 
   auto* a_func =
       Func("a_func", func_params, ty.void_(),
diff --git a/src/writer/spirv/builder_constructor_expression_test.cc b/src/writer/spirv/builder_constructor_expression_test.cc
index cd3cc2f..54fd44c 100644
--- a/src/writer/spirv/builder_constructor_expression_test.cc
+++ b/src/writer/spirv/builder_constructor_expression_test.cc
@@ -138,7 +138,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_IdentifierExpression_Param) {
-  auto* var = Global("ident", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Global("ident", ty.f32(), ast::StorageClass::kFunction);
 
   auto* t = vec2<f32>(1.0f, "ident");
   WrapInFunction(t);
@@ -1285,7 +1285,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_U32_to_I32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<u32>());
+  auto* var = Global("i", ty.vec3<u32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<i32>("i");
   WrapInFunction(cast);
@@ -1311,7 +1311,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_F32_to_I32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* var = Global("i", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<i32>("i");
   WrapInFunction(cast);
@@ -1337,7 +1337,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_I32_to_U32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<i32>());
+  auto* var = Global("i", ty.vec3<i32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<u32>("i");
   WrapInFunction(cast);
@@ -1363,7 +1363,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_F32_to_U32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* var = Global("i", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<u32>("i");
   WrapInFunction(cast);
@@ -1389,7 +1389,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_I32_to_F32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<i32>());
+  auto* var = Global("i", ty.vec3<i32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<f32>("i");
   WrapInFunction(cast);
@@ -1415,7 +1415,7 @@
 }
 
 TEST_F(SpvBuilderConstructorTest, Type_Convert_Vectors_U32_to_F32) {
-  auto* var = Global("i", ast::StorageClass::kPrivate, ty.vec3<u32>());
+  auto* var = Global("i", ty.vec3<u32>(), ast::StorageClass::kPrivate);
 
   auto* cast = vec3<f32>("i");
   WrapInFunction(cast);
@@ -1455,9 +1455,9 @@
 TEST_F(SpvBuilderConstructorTest, IsConstructorConst_GlobalVector_WithIdent) {
   // vec3<f32>(a, b, c)  -> false -- ERROR
 
-  Global("a", ast::StorageClass::kPrivate, ty.f32());
-  Global("b", ast::StorageClass::kPrivate, ty.f32());
-  Global("c", ast::StorageClass::kPrivate, ty.f32());
+  Global("a", ty.f32(), ast::StorageClass::kPrivate);
+  Global("b", ty.f32(), ast::StorageClass::kPrivate);
+  Global("c", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* t = vec3<f32>("a", "b", "c");
   WrapInFunction(t);
@@ -1528,9 +1528,9 @@
   auto* t = vec3<f32>("a", "b", "c");
   WrapInFunction(t);
 
-  Global("a", ast::StorageClass::kPrivate, ty.f32());
-  Global("b", ast::StorageClass::kPrivate, ty.f32());
-  Global("c", ast::StorageClass::kPrivate, ty.f32());
+  Global("a", ty.f32(), ast::StorageClass::kPrivate);
+  Global("b", ty.f32(), ast::StorageClass::kPrivate);
+  Global("c", ty.f32(), ast::StorageClass::kPrivate);
 
   spirv::Builder& b = Build();
 
@@ -1620,8 +1620,8 @@
   auto* t = Construct(s_type, 2.f, "a", 2.f);
   WrapInFunction(t);
 
-  Global("a", ast::StorageClass::kPrivate, ty.f32());
-  Global("b", ast::StorageClass::kPrivate, ty.f32());
+  Global("a", ty.f32(), ast::StorageClass::kPrivate);
+  Global("b", ty.f32(), ast::StorageClass::kPrivate);
 
   spirv::Builder& b = Build();
 
diff --git a/src/writer/spirv/builder_function_decoration_test.cc b/src/writer/spirv/builder_function_decoration_test.cc
index a5f007f..13b8faf 100644
--- a/src/writer/spirv/builder_function_decoration_test.cc
+++ b/src/writer/spirv/builder_function_decoration_test.cc
@@ -97,9 +97,9 @@
                create<ast::StageDecoration>(ast::PipelineStage::kVertex),
            });
 
-  auto* v_in = Global("my_in", ast::StorageClass::kInput, ty.f32());
-  auto* v_out = Global("my_out", ast::StorageClass::kOutput, ty.f32());
-  auto* v_wg = Global("my_wg", ast::StorageClass::kWorkgroup, ty.f32());
+  auto* v_in = Global("my_in", ty.f32(), ast::StorageClass::kInput);
+  auto* v_out = Global("my_out", ty.f32(), ast::StorageClass::kOutput);
+  auto* v_wg = Global("my_wg", ty.f32(), ast::StorageClass::kWorkgroup);
 
   spirv::Builder& b = Build();
 
@@ -142,9 +142,9 @@
                create<ast::StageDecoration>(ast::PipelineStage::kVertex),
            });
 
-  auto* v_in = Global("my_in", ast::StorageClass::kInput, ty.f32());
-  auto* v_out = Global("my_out", ast::StorageClass::kOutput, ty.f32());
-  auto* v_wg = Global("my_wg", ast::StorageClass::kWorkgroup, ty.f32());
+  auto* v_in = Global("my_in", ty.f32(), ast::StorageClass::kInput);
+  auto* v_out = Global("my_out", ty.f32(), ast::StorageClass::kOutput);
+  auto* v_wg = Global("my_wg", ty.f32(), ast::StorageClass::kWorkgroup);
 
   spirv::Builder& b = Build();
 
@@ -258,7 +258,7 @@
 }
 
 TEST_F(BuilderTest, FunctionDecoration_ExecutionMode_FragDepth) {
-  Global("fragdepth", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  Global("fragdepth", ty.f32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kFragDepth),
          });
diff --git a/src/writer/spirv/builder_function_test.cc b/src/writer/spirv/builder_function_test.cc
index 90cadca..807dba3 100644
--- a/src/writer/spirv/builder_function_test.cc
+++ b/src/writer/spirv/builder_function_test.cc
@@ -85,7 +85,7 @@
 }
 
 TEST_F(BuilderTest, Function_Terminator_ReturnValue) {
-  Global("a", ast::StorageClass::kPrivate, ty.f32());
+  Global("a", ty.f32(), ast::StorageClass::kPrivate);
 
   Func("a_func", {}, ty.void_(),
        ast::StatementList{create<ast::ReturnStatement>(Expr("a"))},
@@ -136,8 +136,8 @@
 }
 
 TEST_F(BuilderTest, Function_WithParams) {
-  ast::VariableList params = {Var("a", ast::StorageClass::kFunction, ty.f32()),
-                              Var("b", ast::StorageClass::kFunction, ty.i32())};
+  ast::VariableList params = {Var("a", ty.f32(), ast::StorageClass::kFunction),
+                              Var("b", ty.i32(), ast::StorageClass::kFunction)};
 
   Func("a_func", params, ty.f32(),
        ast::StatementList{create<ast::ReturnStatement>(Expr("a"))},
@@ -238,7 +238,7 @@
   auto* s = ty.struct_("Data", str);
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
 
-  Global("data", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("data", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(0),
@@ -247,7 +247,7 @@
   AST().AddConstructedType(s);
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("a", ast::VariableList{}, ty.void_(),
@@ -261,7 +261,7 @@
   }
 
   {
-    auto* var = Var("v", ast::StorageClass::kFunction, ty.f32(),
+    auto* var = Var("v", ty.f32(), ast::StorageClass::kFunction,
                     MemberAccessor("data", "d"), ast::VariableDecorationList{});
 
     Func("b", ast::VariableList{}, ty.void_(),
diff --git a/src/writer/spirv/builder_function_variable_test.cc b/src/writer/spirv/builder_function_variable_test.cc
index d45f33f..88c9260 100644
--- a/src/writer/spirv/builder_function_variable_test.cc
+++ b/src/writer/spirv/builder_function_variable_test.cc
@@ -45,7 +45,7 @@
 using BuilderTest = TestHelper;
 
 TEST_F(BuilderTest, FunctionVar_NoStorageClass) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -67,7 +67,7 @@
 TEST_F(BuilderTest, FunctionVar_WithConstantConstructor) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput, init,
                    ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -97,7 +97,7 @@
 TEST_F(BuilderTest, FunctionVar_WithNonConstantConstructor) {
   auto* init = vec2<f32>(1.f, Add(3.f, 3.f));
 
-  auto* v = Global("var", ast::StorageClass::kFunction, ty.vec2<f32>(), init,
+  auto* v = Global("var", ty.vec2<f32>(), ast::StorageClass::kFunction, init,
                    ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -129,10 +129,10 @@
   // var v : f32 = 1.0;
   // var v2 : f32 = v; // Should generate the load and store automatically.
 
-  auto* v = Global("v", ast::StorageClass::kFunction, ty.f32(), Expr(1.f),
+  auto* v = Global("v", ty.f32(), ast::StorageClass::kFunction, Expr(1.f),
                    ast::VariableDecorationList{});
 
-  auto* v2 = Global("v2", ast::StorageClass::kFunction, ty.f32(), Expr("v"),
+  auto* v2 = Global("v2", ty.f32(), ast::StorageClass::kFunction, Expr("v"),
                     ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -165,10 +165,10 @@
   // var v : f32 = 1.0;
   // const v2 : f32 = v; // Should generate the load
 
-  auto* v = Global("v", ast::StorageClass::kFunction, ty.f32(), Expr(1.f),
+  auto* v = Global("v", ty.f32(), ast::StorageClass::kFunction, Expr(1.f),
                    ast::VariableDecorationList{});
 
-  auto* v2 = Global("v2", ast::StorageClass::kFunction, ty.f32(), Expr("v"),
+  auto* v2 = Global("v2", ty.f32(), ast::StorageClass::kFunction, Expr("v"),
                     ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -200,7 +200,7 @@
 TEST_F(BuilderTest, FunctionVar_Const) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = Const("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = Const("var", ty.f32(), ast::StorageClass::kOutput, init,
                   ast::VariableDecorationList{});
 
   WrapInFunction(v);
diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc
index c39c281..46c606f 100644
--- a/src/writer/spirv/builder_global_variable_test.cc
+++ b/src/writer/spirv/builder_global_variable_test.cc
@@ -51,7 +51,7 @@
 using BuilderTest = TestHelper;
 
 TEST_F(BuilderTest, GlobalVar_NoStorageClass) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -66,7 +66,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_WithStorageClass) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput);
 
   spirv::Builder& b = Build();
 
@@ -81,7 +81,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_WithStorageClass_Input) {
-  auto* v = Global("var", ast::StorageClass::kInput, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kInput);
 
   spirv::Builder& b = Build();
 
@@ -97,7 +97,7 @@
 TEST_F(BuilderTest, GlobalVar_WithConstructor) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput, init,
                    ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -120,7 +120,7 @@
 TEST_F(BuilderTest, GlobalVar_Const) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = GlobalConst("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = GlobalConst("var", ty.f32(), ast::StorageClass::kOutput, init,
                         ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -141,7 +141,7 @@
 TEST_F(BuilderTest, GlobalVar_Complex_Constructor) {
   auto* init = vec3<f32>(ast::ExpressionList{Expr(1.f), Expr(2.f), Expr(3.f)});
 
-  auto* v = GlobalConst("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = GlobalConst("var", ty.f32(), ast::StorageClass::kOutput, init,
                         ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -161,7 +161,7 @@
 TEST_F(BuilderTest, GlobalVar_Complex_ConstructorWithExtract) {
   auto* init = vec3<f32>(vec2<f32>(1.f, 2.f), 3.f);
 
-  auto* v = GlobalConst("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = GlobalConst("var", ty.f32(), ast::StorageClass::kOutput, init,
                         ast::VariableDecorationList{});
 
   spirv::Builder& b = Build();
@@ -186,7 +186,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_WithLocation) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput, nullptr,
                    ast::VariableDecorationList{
                        create<ast::LocationDecoration>(5),
                    });
@@ -206,7 +206,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_WithBindingAndGroup) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput, nullptr,
                    ast::VariableDecorationList{
                        create<ast::BindingDecoration>(2),
                        create<ast::GroupDecoration>(3),
@@ -228,7 +228,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_WithBuiltin) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32(), nullptr,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput, nullptr,
                    ast::VariableDecorationList{
                        create<ast::BuiltinDecoration>(ast::Builtin::kPosition),
                    });
@@ -248,7 +248,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Bool) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.bool_(), Expr(true),
+  auto* v = Global("var", ty.bool_(), ast::StorageClass::kNone, Expr(true),
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(1200),
                    });
@@ -268,7 +268,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Bool_NoConstructor) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.bool_(), nullptr,
+  auto* v = Global("var", ty.bool_(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(1200),
                    });
@@ -288,7 +288,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Scalar) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.f32(), Expr(2.f),
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kNone, Expr(2.f),
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(0),
                    });
@@ -308,7 +308,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Scalar_F32_NoConstructor) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.f32(), nullptr,
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(0),
                    });
@@ -328,7 +328,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Scalar_I32_NoConstructor) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.i32(), nullptr,
+  auto* v = Global("var", ty.i32(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(0),
                    });
@@ -348,7 +348,7 @@
 }
 
 TEST_F(BuilderTest, GlobalVar_ConstantId_Scalar_U32_NoConstructor) {
-  auto* v = Global("var", ast::StorageClass::kNone, ty.u32(), nullptr,
+  auto* v = Global("var", ty.u32(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::ConstantIdDecoration>(0),
                    });
@@ -419,7 +419,7 @@
                                ast::StructDecorationList{}));
   auto* ac = create<type::AccessControl>(ast::AccessControl::kReadOnly, A);
 
-  auto* var = Global("b", ast::StorageClass::kStorage, ac);
+  auto* var = Global("b", ac, ast::StorageClass::kStorage);
 
   spirv::Builder& b = Build();
 
@@ -452,7 +452,7 @@
                                ast::StructDecorationList{}));
   auto* B = ty.alias("B", A);
   auto* ac = create<type::AccessControl>(ast::AccessControl::kReadOnly, B);
-  auto* var = Global("b", ast::StorageClass::kStorage, ac);
+  auto* var = Global("b", ac, ast::StorageClass::kStorage);
 
   spirv::Builder& b = Build();
 
@@ -483,7 +483,7 @@
                                ast::StructDecorationList{}));
   auto* ac = create<type::AccessControl>(ast::AccessControl::kReadOnly, A);
   auto* B = ty.alias("B", ac);
-  auto* var = Global("b", ast::StorageClass::kStorage, B);
+  auto* var = Global("b", B, ast::StorageClass::kStorage);
 
   spirv::Builder& b = Build();
 
@@ -515,8 +515,8 @@
   type::AccessControl read{ast::AccessControl::kReadOnly, A};
   type::AccessControl rw{ast::AccessControl::kReadWrite, A};
 
-  auto* var_b = Global("b", ast::StorageClass::kStorage, &read);
-  auto* var_c = Global("c", ast::StorageClass::kStorage, &rw);
+  auto* var_b = Global("b", &read, ast::StorageClass::kStorage);
+  auto* var_c = Global("c", &rw, ast::StorageClass::kStorage);
 
   spirv::Builder& b = Build();
 
@@ -553,7 +553,7 @@
 
   auto* ac = create<type::AccessControl>(ast::AccessControl::kReadOnly, type);
 
-  auto* var_a = Global("a", ast::StorageClass::kUniformConstant, ac);
+  auto* var_a = Global("a", ac, ast::StorageClass::kUniformConstant);
 
   spirv::Builder& b = Build();
 
@@ -575,11 +575,11 @@
       type::StorageTexture::SubtypeFor(type::ImageFormat::kR32Uint, Types());
   auto* type = create<type::StorageTexture>(
       type::TextureDimension::k2d, type::ImageFormat::kR32Uint, subtype);
-  Global("test_var", ast::StorageClass::kNone, type);
+  Global("test_var", type, ast::StorageClass::kNone);
 
   auto* ac = create<type::AccessControl>(ast::AccessControl::kWriteOnly, type);
 
-  auto* var_a = Global("a", ast::StorageClass::kUniformConstant, ac);
+  auto* var_a = Global("a", ac, ast::StorageClass::kUniformConstant);
 
   spirv::Builder& b = Build();
 
@@ -605,14 +605,14 @@
   auto* st = create<type::StorageTexture>(type::TextureDimension::k2d,
                                           type::ImageFormat::kR32Uint, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, st);
+  Global("test_var", st, ast::StorageClass::kNone);
 
   auto* type_a = create<type::AccessControl>(ast::AccessControl::kReadOnly, st);
-  auto* var_a = Global("a", ast::StorageClass::kUniformConstant, type_a);
+  auto* var_a = Global("a", type_a, ast::StorageClass::kUniformConstant);
 
   auto* type_b =
       create<type::AccessControl>(ast::AccessControl::kWriteOnly, st);
-  auto* var_b = Global("b", ast::StorageClass::kUniformConstant, type_b);
+  auto* var_b = Global("b", type_b, ast::StorageClass::kUniformConstant);
 
   spirv::Builder& b = Build();
 
@@ -635,7 +635,7 @@
 
 TEST_F(BuilderTest, SampleIndex) {
   auto* var =
-      Global("sample_index", ast::StorageClass::kInput, ty.u32(), nullptr,
+      Global("sample_index", ty.u32(), ast::StorageClass::kInput, nullptr,
              ast::VariableDecorationList{
                  create<ast::BuiltinDecoration>(ast::Builtin::kSampleIndex),
              });
@@ -669,11 +669,11 @@
   //   mask_out[0] = mask_in[0];
   // }
 
-  Global("mask_in", ast::StorageClass::kInput, ty.u32(), nullptr,
+  Global("mask_in", ty.u32(), ast::StorageClass::kInput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kSampleMaskIn),
          });
-  Global("mask_out", ast::StorageClass::kOutput, ty.u32(), nullptr,
+  Global("mask_out", ty.u32(), ast::StorageClass::kOutput, nullptr,
          ast::VariableDecorationList{
              create<ast::BuiltinDecoration>(ast::Builtin::kSampleMaskOut),
          });
diff --git a/src/writer/spirv/builder_ident_expression_test.cc b/src/writer/spirv/builder_ident_expression_test.cc
index 41f289a..89ae49a 100644
--- a/src/writer/spirv/builder_ident_expression_test.cc
+++ b/src/writer/spirv/builder_ident_expression_test.cc
@@ -40,7 +40,7 @@
 TEST_F(BuilderTest, IdentifierExpression_GlobalConst) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = GlobalConst("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = GlobalConst("var", ty.f32(), ast::StorageClass::kOutput, init,
                         ast::VariableDecorationList{});
 
   auto* expr = Expr("var");
@@ -62,7 +62,7 @@
 }
 
 TEST_F(BuilderTest, IdentifierExpression_GlobalVar) {
-  auto* v = Global("var", ast::StorageClass::kOutput, ty.f32());
+  auto* v = Global("var", ty.f32(), ast::StorageClass::kOutput);
 
   auto* expr = Expr("var");
   WrapInFunction(expr);
@@ -85,7 +85,7 @@
 TEST_F(BuilderTest, IdentifierExpression_FunctionConst) {
   auto* init = vec3<f32>(1.f, 1.f, 3.f);
 
-  auto* v = Const("var", ast::StorageClass::kOutput, ty.f32(), init,
+  auto* v = Const("var", ty.f32(), ast::StorageClass::kOutput, init,
                   ast::VariableDecorationList{});
 
   auto* expr = Expr("var");
@@ -107,7 +107,7 @@
 }
 
 TEST_F(BuilderTest, IdentifierExpression_FunctionVar) {
-  auto* v = Var("var", ast::StorageClass::kNone, ty.f32());
+  auto* v = Var("var", ty.f32(), ast::StorageClass::kNone);
   auto* expr = Expr("var");
   WrapInFunction(v, expr);
 
@@ -131,7 +131,7 @@
 }
 
 TEST_F(BuilderTest, IdentifierExpression_Load) {
-  auto* var = Global("var", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("var", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* expr = Add("var", "var");
   WrapInFunction(expr);
@@ -156,7 +156,7 @@
 }
 
 TEST_F(BuilderTest, IdentifierExpression_NoLoadConst) {
-  auto* var = GlobalConst("var", ast::StorageClass::kNone, ty.i32(), Expr(2),
+  auto* var = GlobalConst("var", ty.i32(), ast::StorageClass::kNone, Expr(2),
                           ast::VariableDecorationList{});
 
   auto* expr = Add("var", "var");
diff --git a/src/writer/spirv/builder_if_test.cc b/src/writer/spirv/builder_if_test.cc
index 1921da6..b1e4c3f 100644
--- a/src/writer/spirv/builder_if_test.cc
+++ b/src/writer/spirv/builder_if_test.cc
@@ -92,7 +92,7 @@
   //   v = 2;
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
   auto* expr =
@@ -130,7 +130,7 @@
   //   v = 3;
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
   auto* else_body = create<ast::BlockStatement>(
@@ -176,7 +176,7 @@
   //   v = 3;
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
   auto* else_body = create<ast::BlockStatement>(
@@ -233,7 +233,7 @@
   //   v = 5;
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
   auto* elseif_1_body = create<ast::BlockStatement>(
@@ -560,7 +560,7 @@
   // if (a) {
   // }
 
-  auto* var = Global("a", ast::StorageClass::kFunction, ty.bool_());
+  auto* var = Global("a", ty.bool_(), ast::StorageClass::kFunction);
 
   auto* expr = create<ast::IfStatement>(
       Expr("a"), create<ast::BlockStatement>(ast::StatementList{}),
diff --git a/src/writer/spirv/builder_intrinsic_test.cc b/src/writer/spirv/builder_intrinsic_test.cc
index b20e553..3f5e17c 100644
--- a/src/writer/spirv/builder_intrinsic_test.cc
+++ b/src/writer/spirv/builder_intrinsic_test.cc
@@ -72,7 +72,7 @@
 TEST_P(IntrinsicBoolTest, Call_Bool) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<bool>());
+  auto* var = Global("v", ty.vec3<bool>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -103,7 +103,7 @@
 TEST_P(IntrinsicFloatTest, Call_Float_Scalar) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("v", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -129,7 +129,7 @@
 TEST_P(IntrinsicFloatTest, Call_Float_Vector) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -162,7 +162,7 @@
 TEST_P(IntrinsicIntTest, Call_SInt_Scalar) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -187,7 +187,7 @@
 TEST_P(IntrinsicIntTest, Call_SInt_Vector) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<i32>());
+  auto* var = Global("v", ty.vec3<i32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -213,7 +213,7 @@
 TEST_P(IntrinsicIntTest, Call_UInt_Scalar) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.u32());
+  auto* var = Global("v", ty.u32(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -238,7 +238,7 @@
 TEST_P(IntrinsicIntTest, Call_UInt_Vector) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<u32>());
+  auto* var = Global("v", ty.vec3<u32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -267,7 +267,7 @@
                     IntrinsicData{"reverseBits", "OpBitReverse"}));
 
 TEST_F(IntrinsicBuilderTest, Call_Dot) {
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("dot", "v", "v");
   WrapInFunction(expr);
@@ -295,7 +295,7 @@
 TEST_P(IntrinsicDeriveTest, Call_Derivative_Scalar) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("v", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -320,7 +320,7 @@
 TEST_P(IntrinsicDeriveTest, Call_Derivative_Vector) {
   auto param = GetParam();
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* var = Global("v", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call(param.name, "v");
   WrapInFunction(expr);
@@ -363,10 +363,10 @@
                     IntrinsicData{"fwidthCoarse", "OpFwidthCoarse"}));
 
 TEST_F(IntrinsicBuilderTest, Call_Select) {
-  auto* v3 = Global("v3", ast::StorageClass::kPrivate, ty.vec3<f32>());
+  auto* v3 = Global("v3", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
   auto* bool_v3 =
-      Global("bool_v3", ast::StorageClass::kPrivate, ty.vec3<bool>());
+      Global("bool_v3", ty.vec3<bool>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("select", "v3", "v3", "bool_v3");
   WrapInFunction(expr);
@@ -403,9 +403,9 @@
   type::Sampler s(type::SamplerKind::kComparisonSampler);
   type::DepthTexture t(type::TextureDimension::k2d);
 
-  auto* tex = Global("texture", ast::StorageClass::kNone, &t);
+  auto* tex = Global("texture", &t, ast::StorageClass::kNone);
 
-  auto* sampler = Global("sampler", ast::StorageClass::kNone, &s);
+  auto* sampler = Global("sampler", &s, ast::StorageClass::kNone);
 
   auto* expr1 = Call("textureSampleCompare", "texture", "sampler",
                      vec2<f32>(1.0f, 2.0f), 2.0f);
@@ -453,7 +453,7 @@
 }
 
 TEST_F(IntrinsicBuilderTest, Call_GLSLMethod_WithLoad) {
-  auto* var = Global("ident", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("ident", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("round", "ident");
   WrapInFunction(expr);
@@ -1283,7 +1283,7 @@
                          testing::Values(IntrinsicData{"clamp", "UClamp"}));
 
 TEST_F(IntrinsicBuilderTest, Call_Modf) {
-  auto* out = Var("out", ast::StorageClass::kFunction, ty.vec2<f32>());
+  auto* out = Var("out", ty.vec2<f32>(), ast::StorageClass::kFunction);
   auto* expr = Call("modf", vec2<f32>(1.0f, 2.0f), "out");
   Func("a_func", ast::VariableList{}, ty.void_(),
        ast::StatementList{
@@ -1326,7 +1326,7 @@
 }
 
 TEST_F(IntrinsicBuilderTest, Call_Frexp) {
-  auto* out = Var("out", ast::StorageClass::kFunction, ty.vec2<i32>());
+  auto* out = Var("out", ty.vec2<i32>(), ast::StorageClass::kFunction);
   auto* expr = Call("frexp", vec2<f32>(1.0f, 2.0f), "out");
   Func("a_func", ast::VariableList{}, ty.void_(),
        ast::StatementList{
@@ -1371,7 +1371,7 @@
 }
 
 TEST_F(IntrinsicBuilderTest, Call_Determinant) {
-  auto* var = Global("var", ast::StorageClass::kPrivate, ty.mat3x3<f32>());
+  auto* var = Global("var", ty.mat3x3<f32>(), ast::StorageClass::kPrivate);
 
   auto* expr = Call("determinant", "var");
   WrapInFunction(expr);
@@ -1413,7 +1413,7 @@
           create<ast::StructBlockDecoration>(),
       });
   auto* s_type = ty.struct_("my_struct", s);
-  Global("b", ast::StorageClass::kStorage, s_type, nullptr,
+  Global("b", s_type, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(1),
              create<ast::GroupDecoration>(2),
@@ -1464,7 +1464,7 @@
       });
 
   auto* s_type = ty.struct_("my_struct", s);
-  Global("b", ast::StorageClass::kStorage, s_type, nullptr,
+  Global("b", s_type, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(1),
              create<ast::GroupDecoration>(2),
diff --git a/src/writer/spirv/builder_loop_test.cc b/src/writer/spirv/builder_loop_test.cc
index 9da5ad6..48ef0c0 100644
--- a/src/writer/spirv/builder_loop_test.cc
+++ b/src/writer/spirv/builder_loop_test.cc
@@ -67,7 +67,7 @@
   //   v = 2;
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
 
@@ -109,7 +109,7 @@
   //   }
   // }
 
-  auto* var = Global("v", ast::StorageClass::kPrivate, ty.i32());
+  auto* var = Global("v", ty.i32(), ast::StorageClass::kPrivate);
   auto* body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(2))});
   auto* continuing = create<ast::BlockStatement>(
diff --git a/src/writer/spirv/builder_return_test.cc b/src/writer/spirv/builder_return_test.cc
index 0227b03..689769d 100644
--- a/src/writer/spirv/builder_return_test.cc
+++ b/src/writer/spirv/builder_return_test.cc
@@ -71,7 +71,7 @@
 }
 
 TEST_F(BuilderTest, Return_WithValue_GeneratesLoad) {
-  auto* var = Global("param", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Global("param", ty.f32(), ast::StorageClass::kFunction);
 
   auto* ret = create<ast::ReturnStatement>(Expr("param"));
   WrapInFunction(ret);
diff --git a/src/writer/spirv/builder_switch_test.cc b/src/writer/spirv/builder_switch_test.cc
index 219afd8..c0ee18f 100644
--- a/src/writer/spirv/builder_switch_test.cc
+++ b/src/writer/spirv/builder_switch_test.cc
@@ -72,8 +72,8 @@
   //     v = 2;
   // }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* case_1_body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(1))});
@@ -142,8 +142,8 @@
   //     v = 1;
   //  }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* default_body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(1))});
@@ -201,8 +201,8 @@
   //      v = 3;
   //  }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* case_1_body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(1))});
@@ -284,8 +284,8 @@
   //      v = 3;
   //  }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* case_1_body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(1)),
@@ -363,8 +363,8 @@
   //      fallthrough;
   //  }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* case_1_body = create<ast::BlockStatement>(
       ast::StatementList{create<ast::AssignmentStatement>(Expr("v"), Expr(1)),
@@ -402,8 +402,8 @@
   //     v = 1;
   //  }
 
-  auto* v = Global("v", ast::StorageClass::kPrivate, ty.i32());
-  auto* a = Global("a", ast::StorageClass::kPrivate, ty.i32());
+  auto* v = Global("v", ty.i32(), ast::StorageClass::kPrivate);
+  auto* a = Global("a", ty.i32(), ast::StorageClass::kPrivate);
 
   auto* if_body = create<ast::BlockStatement>(ast::StatementList{
       create<ast::BreakStatement>(),
diff --git a/src/writer/spirv/builder_type_test.cc b/src/writer/spirv/builder_type_test.cc
index 2ffa217..d9d3f3e 100644
--- a/src/writer/spirv/builder_type_test.cc
+++ b/src/writer/spirv/builder_type_test.cc
@@ -879,7 +879,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1d,
                                          type::ImageFormat::kR16Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -901,7 +901,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1d,
                                          type::ImageFormat::kR8Snorm, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -923,7 +923,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1d,
                                          type::ImageFormat::kR8Unorm, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -945,7 +945,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1d,
                                          type::ImageFormat::kR8Uint, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -962,7 +962,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1d,
                                          type::ImageFormat::kR8Sint, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -979,7 +979,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k1dArray,
                                          type::ImageFormat::kR16Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1001,7 +1001,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k2d,
                                          type::ImageFormat::kR16Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1018,7 +1018,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k2dArray,
                                          type::ImageFormat::kR16Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1035,7 +1035,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k3d,
                                          type::ImageFormat::kR16Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1053,7 +1053,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k2d,
                                          type::ImageFormat::kR32Float, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1071,7 +1071,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k2d,
                                          type::ImageFormat::kR32Sint, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
@@ -1089,7 +1089,7 @@
   auto* s = create<type::StorageTexture>(type::TextureDimension::k2d,
                                          type::ImageFormat::kR32Uint, subtype);
 
-  Global("test_var", ast::StorageClass::kNone, s);
+  Global("test_var", s, ast::StorageClass::kNone);
 
   spirv::Builder& b = Build();
 
diff --git a/src/writer/spirv/builder_unary_op_expression_test.cc b/src/writer/spirv/builder_unary_op_expression_test.cc
index 0eb0d1e..4034dbf 100644
--- a/src/writer/spirv/builder_unary_op_expression_test.cc
+++ b/src/writer/spirv/builder_unary_op_expression_test.cc
@@ -87,7 +87,7 @@
 }
 
 TEST_F(BuilderTest, UnaryOp_LoadRequired) {
-  auto* var = Global("param", ast::StorageClass::kFunction, ty.vec3<f32>());
+  auto* var = Global("param", ty.vec3<f32>(), ast::StorageClass::kFunction);
 
   auto* expr =
       create<ast::UnaryOpExpression>(ast::UnaryOp::kNegation, Expr("param"));
diff --git a/src/writer/wgsl/generator_impl_entry_point_test.cc b/src/writer/wgsl/generator_impl_entry_point_test.cc
index 057629d..7f98d0e 100644
--- a/src/writer/wgsl/generator_impl_entry_point_test.cc
+++ b/src/writer/wgsl/generator_impl_entry_point_test.cc
@@ -65,11 +65,11 @@
 
 TEST_F(WgslGeneratorImplTest, Emit_EntryPoint_UnusedVariable) {
   auto* global_unused =
-      Global("global_unused", ast::StorageClass::kInput, ty.f32());
+      Global("global_unused", ty.f32(), ast::StorageClass::kInput);
   create<ast::VariableDeclStatement>(global_unused);
 
   auto* global_used =
-      Global("global_used", ast::StorageClass::kInput, ty.f32());
+      Global("global_used", ty.f32(), ast::StorageClass::kInput);
   create<ast::VariableDeclStatement>(global_used);
 
   Func("main", ast::VariableList{}, ty.void_(),
@@ -97,7 +97,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, Emit_EntryPoint_GlobalsInterleaved) {
-  auto* global0 = Global("a0", ast::StorageClass::kInput, ty.f32());
+  auto* global0 = Global("a0", ty.f32(), ast::StorageClass::kInput);
   create<ast::VariableDeclStatement>(global0);
 
   auto* str0 = create<ast::Struct>(ast::StructMemberList{Member("a", ty.i32())},
@@ -111,7 +111,7 @@
        },
        ast::FunctionDecorationList{});
 
-  auto* global1 = Global("a1", ast::StorageClass::kOutput, ty.f32());
+  auto* global1 = Global("a1", ty.f32(), ast::StorageClass::kOutput);
   create<ast::VariableDeclStatement>(global1);
 
   auto* str1 = create<ast::Struct>(ast::StructMemberList{Member("a", ty.i32())},
@@ -124,9 +124,9 @@
   Func("main", ast::VariableList{}, ty.void_(),
        ast::StatementList{
            create<ast::VariableDeclStatement>(
-               Var("s0", ast::StorageClass::kFunction, s0)),
+               Var("s0", s0, ast::StorageClass::kFunction)),
            create<ast::VariableDeclStatement>(
-               Var("s1", ast::StorageClass::kFunction, s1)),
+               Var("s1", s1, ast::StorageClass::kFunction)),
            create<ast::AssignmentStatement>(Expr("a1"), Expr(call_func)),
        },
        ast::FunctionDecorationList{
diff --git a/src/writer/wgsl/generator_impl_function_test.cc b/src/writer/wgsl/generator_impl_function_test.cc
index 9414cde..f99f512 100644
--- a/src/writer/wgsl/generator_impl_function_test.cc
+++ b/src/writer/wgsl/generator_impl_function_test.cc
@@ -63,8 +63,8 @@
 TEST_F(WgslGeneratorImplTest, Emit_Function_WithParams) {
   auto* func =
       Func("my_func",
-           ast::VariableList{Var("a", ast::StorageClass::kNone, ty.f32()),
-                             Var("b", ast::StorageClass::kNone, ty.i32())},
+           ast::VariableList{Var("a", ty.f32(), ast::StorageClass::kNone),
+                             Var("b", ty.i32(), ast::StorageClass::kNone)},
            ty.void_(),
            ast::StatementList{
                create<ast::DiscardStatement>(),
@@ -185,7 +185,7 @@
   type::AccessControl ac(ast::AccessControl::kReadWrite, s);
   AST().AddConstructedType(s);
 
-  Global("data", ast::StorageClass::kStorage, &ac, nullptr,
+  Global("data", &ac, ast::StorageClass::kStorage, nullptr,
          ast::VariableDecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(0),
@@ -193,7 +193,7 @@
 
   {
     auto* var =
-        Var("v", ast::StorageClass::kFunction, ty.f32(),
+        Var("v", ty.f32(), ast::StorageClass::kFunction,
             create<ast::MemberAccessorExpression>(Expr("data"), Expr("d")),
             ast::VariableDecorationList{});
 
@@ -209,7 +209,7 @@
 
   {
     auto* var =
-        Var("v", ast::StorageClass::kFunction, ty.f32(),
+        Var("v", ty.f32(), ast::StorageClass::kFunction,
             create<ast::MemberAccessorExpression>(Expr("data"), Expr("d")),
             ast::VariableDecorationList{});
 
diff --git a/src/writer/wgsl/generator_impl_global_decl_test.cc b/src/writer/wgsl/generator_impl_global_decl_test.cc
index 464ee70..9f72bc2 100644
--- a/src/writer/wgsl/generator_impl_global_decl_test.cc
+++ b/src/writer/wgsl/generator_impl_global_decl_test.cc
@@ -30,11 +30,11 @@
 using WgslGeneratorImplTest = TestHelper;
 
 TEST_F(WgslGeneratorImplTest, Emit_GlobalDeclAfterFunction) {
-  auto* func_var = Var("a", ast::StorageClass::kFunction, ty.f32(), nullptr,
+  auto* func_var = Var("a", ty.f32(), ast::StorageClass::kFunction, nullptr,
                        ast::VariableDecorationList{});
   WrapInFunction(create<ast::VariableDeclStatement>(func_var));
 
-  auto* global_var = Global("a", ast::StorageClass::kInput, ty.f32());
+  auto* global_var = Global("a", ty.f32(), ast::StorageClass::kInput);
   create<ast::VariableDeclStatement>(global_var);
 
   GeneratorImpl& gen = Build();
@@ -51,7 +51,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, Emit_GlobalsInterleaved) {
-  auto* global0 = Global("a0", ast::StorageClass::kInput, ty.f32());
+  auto* global0 = Global("a0", ty.f32(), ast::StorageClass::kInput);
   create<ast::VariableDeclStatement>(global0);
 
   auto* str0 = create<ast::Struct>(ast::StructMemberList{Member("a", ty.i32())},
@@ -65,7 +65,7 @@
        },
        ast::FunctionDecorationList{});
 
-  auto* global1 = Global("a1", ast::StorageClass::kOutput, ty.f32());
+  auto* global1 = Global("a1", ty.f32(), ast::StorageClass::kOutput);
   create<ast::VariableDeclStatement>(global1);
 
   auto* str1 = create<ast::Struct>(ast::StructMemberList{Member("a", ty.i32())},
@@ -78,9 +78,9 @@
   Func("main", ast::VariableList{}, ty.void_(),
        ast::StatementList{
            create<ast::VariableDeclStatement>(
-               Var("s0", ast::StorageClass::kFunction, s0)),
+               Var("s0", s0, ast::StorageClass::kFunction)),
            create<ast::VariableDeclStatement>(
-               Var("s1", ast::StorageClass::kFunction, s1)),
+               Var("s1", s1, ast::StorageClass::kFunction)),
            create<ast::AssignmentStatement>(Expr("a1"), Expr(call_func)),
        },
        ast::FunctionDecorationList{
diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc
index 8b25c28..c1d4c77 100644
--- a/src/writer/wgsl/generator_impl_test.cc
+++ b/src/writer/wgsl/generator_impl_test.cc
@@ -55,7 +55,7 @@
 TEST_P(WgslBuiltinConversionTest, Emit) {
   auto params = GetParam();
 
-  auto* var = Global("a", ast::StorageClass::kNone, ty.f32(), nullptr,
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kNone, nullptr,
                      ast::VariableDecorationList{
                          create<ast::BuiltinDecoration>(params.builtin),
                      });
diff --git a/src/writer/wgsl/generator_impl_variable_decl_statement_test.cc b/src/writer/wgsl/generator_impl_variable_decl_statement_test.cc
index 3290824..ce2d3e2 100644
--- a/src/writer/wgsl/generator_impl_variable_decl_statement_test.cc
+++ b/src/writer/wgsl/generator_impl_variable_decl_statement_test.cc
@@ -33,7 +33,7 @@
 using WgslGeneratorImplTest = TestHelper;
 
 TEST_F(WgslGeneratorImplTest, Emit_VariableDeclStatement) {
-  auto* var = Global("a", ast::StorageClass::kNone, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kNone);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
@@ -51,7 +51,7 @@
   // storage class.  Rely on defaulting.
   // https://github.com/gpuweb/gpuweb/issues/654
 
-  auto* var = Global("a", ast::StorageClass::kFunction, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kFunction);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
@@ -65,7 +65,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, Emit_VariableDeclStatement_Private) {
-  auto* var = Global("a", ast::StorageClass::kPrivate, ty.f32());
+  auto* var = Global("a", ty.f32(), ast::StorageClass::kPrivate);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
   WrapInFunction(stmt);
@@ -79,8 +79,8 @@
 }
 
 TEST_F(WgslGeneratorImplTest, Emit_VariableDeclStatement_Sampler) {
-  auto* var = Global("s", ast::StorageClass::kUniformConstant,
-                     create<type::Sampler>(type::SamplerKind::kSampler));
+  auto* var = Global("s", create<type::Sampler>(type::SamplerKind::kSampler),
+                     ast::StorageClass::kUniformConstant);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
@@ -95,9 +95,9 @@
 TEST_F(WgslGeneratorImplTest, Emit_VariableDeclStatement_Texture) {
   auto* st =
       create<type::SampledTexture>(type::TextureDimension::k1d, ty.f32());
-  auto* var =
-      Global("t", ast::StorageClass::kUniformConstant,
-             create<type::AccessControl>(ast::AccessControl::kReadOnly, st));
+  auto* var = Global(
+      "t", create<type::AccessControl>(ast::AccessControl::kReadOnly, st),
+      ast::StorageClass::kUniformConstant);
 
   auto* stmt = create<ast::VariableDeclStatement>(var);
 
diff --git a/src/writer/wgsl/generator_impl_variable_test.cc b/src/writer/wgsl/generator_impl_variable_test.cc
index 31f85dd..09c3b40 100644
--- a/src/writer/wgsl/generator_impl_variable_test.cc
+++ b/src/writer/wgsl/generator_impl_variable_test.cc
@@ -34,7 +34,7 @@
 using WgslGeneratorImplTest = TestHelper;
 
 TEST_F(WgslGeneratorImplTest, EmitVariable) {
-  auto* v = Global("a", ast::StorageClass::kNone, ty.f32());
+  auto* v = Global("a", ty.f32(), ast::StorageClass::kNone);
 
   GeneratorImpl& gen = Build();
 
@@ -44,7 +44,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_StorageClass) {
-  auto* v = Global("a", ast::StorageClass::kInput, ty.f32());
+  auto* v = Global("a", ty.f32(), ast::StorageClass::kInput);
 
   GeneratorImpl& gen = Build();
 
@@ -54,7 +54,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Decorated) {
-  auto* v = Global("a", ast::StorageClass::kNone, ty.f32(), nullptr,
+  auto* v = Global("a", ty.f32(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::LocationDecoration>(2),
                    });
@@ -67,7 +67,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Decorated_Multiple) {
-  auto* v = Global("a", ast::StorageClass::kNone, ty.f32(), nullptr,
+  auto* v = Global("a", ty.f32(), ast::StorageClass::kNone, nullptr,
                    ast::VariableDecorationList{
                        create<ast::BuiltinDecoration>(ast::Builtin::kPosition),
                        create<ast::BindingDecoration>(0),
@@ -86,7 +86,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Constructor) {
-  auto* v = Global("a", ast::StorageClass::kNone, ty.f32(), Expr("initializer"),
+  auto* v = Global("a", ty.f32(), ast::StorageClass::kNone, Expr("initializer"),
                    ast::VariableDecorationList{});
 
   GeneratorImpl& gen = Build();
@@ -97,7 +97,7 @@
 }
 
 TEST_F(WgslGeneratorImplTest, EmitVariable_Const) {
-  auto* v = Const("a", ast::StorageClass::kNone, ty.f32(), Expr("initializer"),
+  auto* v = Const("a", ty.f32(), ast::StorageClass::kNone, Expr("initializer"),
                   ast::VariableDecorationList{});
 
   GeneratorImpl& gen = Build();