writer/msl: Emit builtins as parameters

Add a config parameter for the CanonicalizeEntryPoint transform that
selects between emitting builtins as parameters (for MSL) or struct
members (for HLSL).

This fixes all of the shader IO issues in Tint's E2E tests for MSL.

Fixed: tint:817
Change-Id: Ieb31cdbd2e4d96ac41f8d8515fd07ead8241d770
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/53282
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>

src/transform/canonicalize_entry_point_io.cc

@@ -24,6 +24,8 @@
 #include "src/sem/struct.h"
 #include "src/sem/variable.h"
 
+TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO::Config);
+
 namespace tint {
 namespace transform {
 
@@ -60,10 +62,17 @@ bool StructMemberComparator(const ast::StructMember* a,
 
 }  // namespace
 
-Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
+Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap& data) {
   ProgramBuilder out;
   CloneContext ctx(&out, in);
 
+  auto* cfg = data.Get<Config>();
+  if (cfg == nullptr) {
+    out.Diagnostics().add_error(
+        "missing transform data for CanonicalizeEntryPointIO");
+    return Output(Program(std::move(out)));
+  }
+
   // Strip entry point IO decorations from struct declarations.
   // TODO(jrprice): This code is duplicated with the SPIR-V transform.
   for (auto* ty : ctx.src->AST().ConstructedTypes()) {
@@ -104,6 +113,16 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
       auto new_struct_param_symbol = ctx.dst->Sym();
       ast::StructMemberList new_struct_members;
       for (auto* param : func->Parameters()) {
+        if (cfg->builtin_style == BuiltinStyle::kParameter &&
+            ast::HasDecoration<ast::BuiltinDecoration>(
+                param->Declaration()->decorations())) {
+          // If this parameter is a builtin and we are emitting those as
+          // parameters, then just clone it as is.
+          new_parameters.push_back(
+              ctx.Clone(const_cast<ast::Variable*>(param->Declaration())));
+          continue;
+        }
+
         auto param_name = ctx.Clone(param->Declaration()->symbol());
         auto* param_ty = param->Type();
         auto* param_declared_ty = param->Declaration()->type();
@@ -118,6 +137,20 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
               TINT_ICE(ctx.dst->Diagnostics()) << "nested pipeline IO struct";
             }
 
+            if (cfg->builtin_style == BuiltinStyle::kParameter &&
+                ast::HasDecoration<ast::BuiltinDecoration>(
+                    member->Declaration()->decorations())) {
+              // If this struct member is a builtin and we are emitting those as
+              // parameters, then move it to the parameter list.
+              auto* member_ty = CreateASTTypeFor(&ctx, member->Type());
+              auto new_param_name = ctx.dst->Sym();
+              new_parameters.push_back(ctx.dst->Param(
+                  new_param_name, member_ty,
+                  ctx.Clone(member->Declaration()->decorations())));
+              init_values.push_back(ctx.dst->Expr(new_param_name));
+              continue;
+            }
+
             ast::DecorationList new_decorations = RemoveDecorations(
                 &ctx, member->Declaration()->decorations(),
                 [](const ast::Decoration* deco) {
@@ -161,6 +194,7 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
         }
       }
 
+      if (!new_struct_members.empty()) {
         // Sort struct members to satisfy HLSL interfacing matching rules.
         std::sort(new_struct_members.begin(), new_struct_members.end(),
                   StructMemberComparator);
@@ -177,6 +211,7 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
             new_struct_param_symbol, ctx.dst->ty.type_name(in_struct_name));
         new_parameters.push_back(struct_param);
       }
+    }
 
     // Handle return type.
     auto* ret_type = func->ReturnType();
@@ -273,5 +308,13 @@ Output CanonicalizeEntryPointIO::Run(const Program* in, const DataMap&) {
   return Output(Program(std::move(out)));
 }
 
+CanonicalizeEntryPointIO::Config::Config(BuiltinStyle builtins)
+    : builtin_style(builtins) {}
+
+CanonicalizeEntryPointIO::Config::Config(const Config&) = default;
+CanonicalizeEntryPointIO::Config::~Config() = default;
+CanonicalizeEntryPointIO::Config& CanonicalizeEntryPointIO::Config::operator=(
+    const Config&) = default;
+
 }  // namespace transform
 }  // namespace tint

src/transform/canonicalize_entry_point_io.h

@@ -70,6 +70,34 @@ namespace transform {
 /// ```
 class CanonicalizeEntryPointIO : public Transform {
  public:
+  /// BuiltinStyle is an enumerator of different ways to emit builtins.
+  enum class BuiltinStyle {
+    /// Use non-struct function parameters for all builtins.
+    kParameter,
+    /// Use struct members for all builtins.
+    kStructMember,
+  };
+
+  /// Configuration options for the transform.
+  struct Config : public Castable<Config, Data> {
+    /// Constructor
+    /// @param builtins the approach to use for emitting builtins.
+    explicit Config(BuiltinStyle builtins);
+
+    /// Copy constructor
+    Config(const Config&);
+
+    /// Destructor
+    ~Config() override;
+
+    /// Assignment operator
+    /// @returns this Config
+    Config& operator=(const Config&);
+
+    /// The approach to use for emitting builtins.
+    BuiltinStyle builtin_style;
+  };
+
   /// Constructor
   CanonicalizeEntryPointIO();
   ~CanonicalizeEntryPointIO() override;

src/transform/canonicalize_entry_point_io_test.cc

@@ -22,7 +22,51 @@ namespace {
 
 using CanonicalizeEntryPointIOTest = TransformTest;
 
-TEST_F(CanonicalizeEntryPointIOTest, Parameters) {
+TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) {
+  auto* src = "";
+
+  auto* expect = "error: missing transform data for CanonicalizeEntryPointIO";
+
+  auto got = Run<CanonicalizeEntryPointIO>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CanonicalizeEntryPointIOTest, Parameters_BuiltinsAsParameters) {
+  auto* src = R"(
+[[stage(fragment)]]
+fn frag_main([[location(1)]] loc1 : f32,
+             [[location(2)]] loc2 : vec4<u32>,
+             [[builtin(position)]] coord : vec4<f32>) {
+  var col : f32 = (coord.x * loc1);
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  [[location(1)]]
+  loc1 : f32;
+  [[location(2)]]
+  loc2 : vec4<u32>;
+};
+
+[[stage(fragment)]]
+fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_1) {
+  let loc1 : f32 = tint_symbol.loc1;
+  let loc2 : vec4<u32> = tint_symbol.loc2;
+  var col : f32 = (coord.x * loc1);
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CanonicalizeEntryPointIOTest, Parameters_BuiltinsAsStructMembers) {
   auto* src = R"(
 [[stage(fragment)]]
 fn frag_main([[location(1)]] loc1 : f32,
@@ -51,7 +95,10 @@ fn frag_main(tint_symbol : tint_symbol_1) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -81,12 +128,49 @@ fn frag_main(tint_symbol : tint_symbol_1) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(CanonicalizeEntryPointIOTest, Parameters_EmptyBody) {
+TEST_F(CanonicalizeEntryPointIOTest,
+       Parameters_EmptyBody_BuiltinsAsParameters) {
+  auto* src = R"(
+[[stage(fragment)]]
+fn frag_main([[location(1)]] loc1 : f32,
+             [[location(2)]] loc2 : vec4<u32>,
+             [[builtin(position)]] coord : vec4<f32>) {
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  [[location(1)]]
+  loc1 : f32;
+  [[location(2)]]
+  loc2 : vec4<u32>;
+};
+
+[[stage(fragment)]]
+fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_1) {
+  let loc1 : f32 = tint_symbol.loc1;
+  let loc2 : vec4<u32> = tint_symbol.loc2;
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CanonicalizeEntryPointIOTest,
+       Parameters_EmptyBody_BuiltinsAsStructMembers) {
   auto* src = R"(
 [[stage(fragment)]]
 fn frag_main([[location(1)]] loc1 : f32,
@@ -113,12 +197,69 @@ fn frag_main(tint_symbol : tint_symbol_1) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(CanonicalizeEntryPointIOTest, StructParameters) {
+TEST_F(CanonicalizeEntryPointIOTest, StructParameters_BuiltinsAsParameters) {
+  auto* src = R"(
+struct FragBuiltins {
+  [[builtin(position)]] coord : vec4<f32>;
+};
+struct FragLocations {
+  [[location(1)]] loc1 : f32;
+  [[location(2)]] loc2 : vec4<u32>;
+};
+
+[[stage(fragment)]]
+fn frag_main([[location(0)]] loc0 : f32,
+             locations : FragLocations,
+             builtins : FragBuiltins) {
+  var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
+}
+)";
+
+  auto* expect = R"(
+struct FragBuiltins {
+  coord : vec4<f32>;
+};
+
+struct FragLocations {
+  loc1 : f32;
+  loc2 : vec4<u32>;
+};
+
+struct tint_symbol_2 {
+  [[location(0)]]
+  loc0 : f32;
+  [[location(1)]]
+  loc1 : f32;
+  [[location(2)]]
+  loc2 : vec4<u32>;
+};
+
+[[stage(fragment)]]
+fn frag_main([[builtin(position)]] tint_symbol_1 : vec4<f32>, tint_symbol : tint_symbol_2) {
+  let loc0 : f32 = tint_symbol.loc0;
+  let locations : FragLocations = FragLocations(tint_symbol.loc1, tint_symbol.loc2);
+  let builtins : FragBuiltins = FragBuiltins(tint_symbol_1);
+  var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CanonicalizeEntryPointIOTest, StructParameters_BuiltinsAsStructMembers) {
   auto* src = R"(
 struct FragBuiltins {
   [[builtin(position)]] coord : vec4<f32>;
@@ -166,7 +307,10 @@ fn frag_main(tint_symbol : tint_symbol_1) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -191,7 +335,10 @@ fn frag_main() -> tint_symbol {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -240,7 +387,10 @@ fn frag_main() -> tint_symbol {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -304,7 +454,10 @@ fn frag_main2(tint_symbol_2 : tint_symbol_3) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -366,7 +519,10 @@ fn frag_main1(tint_symbol : tint_symbol_1) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -444,7 +600,10 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -501,7 +660,10 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -593,7 +755,10 @@ fn frag_main(tint_symbol_2 : tint_symbol_3) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -617,7 +782,10 @@ fn tint_symbol_1(tint_symbol : tint_symbol_2) {
 }
 )";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }

src/transform/hlsl.cc

@@ -32,8 +32,9 @@ namespace transform {
 Hlsl::Hlsl() = default;
 Hlsl::~Hlsl() = default;
 
-Output Hlsl::Run(const Program* in, const DataMap& data) {
+Output Hlsl::Run(const Program* in, const DataMap&) {
   Manager manager;
+  DataMap data;
   manager.Add<CanonicalizeEntryPointIO>();
   manager.Add<InlinePointerLets>();
   // Simplify cleans up messy `*(&(expr))` expressions from InlinePointerLets.
@@ -46,6 +47,8 @@ Output Hlsl::Run(const Program* in, const DataMap& data) {
   manager.Add<CalculateArrayLength>();
   manager.Add<ExternalTextureTransform>();
   manager.Add<PromoteInitializersToConstVar>();
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
   auto out = manager.Run(in, data);
   if (!out.program.IsValid()) {
     return out;

src/transform/msl.cc

@@ -35,11 +35,14 @@ namespace transform {
 Msl::Msl() = default;
 Msl::~Msl() = default;
 
-Output Msl::Run(const Program* in, const DataMap& data) {
+Output Msl::Run(const Program* in, const DataMap&) {
   Manager manager;
+  DataMap data;
   manager.Add<CanonicalizeEntryPointIO>();
   manager.Add<ExternalTextureTransform>();
   manager.Add<PromoteInitializersToConstVar>();
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
   auto out = manager.Run(in, data);
   if (!out.program.IsValid()) {
     return out;

src/writer/msl/generator_impl_function_test.cc

@@ -140,16 +140,12 @@ TEST_F(MslGeneratorImplTest, Emit_Decoration_EntryPoint_WithInOut_Builtins) {
 
 using namespace metal;
 struct tint_symbol_1 {
-  float4 coord [[position]];
-};
-struct tint_symbol_2 {
   float value [[depth(any)]];
 };
 
-fragment tint_symbol_2 frag_main(tint_symbol_1 tint_symbol [[stage_in]]) {
+fragment tint_symbol_1 frag_main(float4 coord [[position]]) {
-  float4 const coord = tint_symbol.coord;
+  tint_symbol_1 const tint_symbol_2 = {.value=coord.x};
-  tint_symbol_2 const tint_symbol_3 = {.value=coord.x};
+  return tint_symbol_2;
-  return tint_symbol_3;
 }
 
 )");
@@ -209,20 +205,19 @@ struct tint_symbol {
   float col2 [[user(locn2)]];
   float4 pos [[position]];
 };
-struct tint_symbol_3 {
+struct tint_symbol_4 {
   float col1 [[user(locn1)]];
   float col2 [[user(locn2)]];
-  float4 pos [[position]];
 };
 
 vertex tint_symbol vert_main() {
   Interface const tint_symbol_1 = {.col1=0.5f, .col2=0.25f, .pos=float4()};
-  tint_symbol const tint_symbol_4 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=tint_symbol_1.pos};
+  tint_symbol const tint_symbol_5 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=tint_symbol_1.pos};
-  return tint_symbol_4;
+  return tint_symbol_5;
 }
 
-fragment void frag_main(tint_symbol_3 tint_symbol_2 [[stage_in]]) {
+fragment void frag_main(float4 tint_symbol_3 [[position]], tint_symbol_4 tint_symbol_2 [[stage_in]]) {
-  Interface const colors = {.col1=tint_symbol_2.col1, .col2=tint_symbol_2.col2, .pos=tint_symbol_2.pos};
+  Interface const colors = {.col1=tint_symbol_2.col1, .col2=tint_symbol_2.col2, .pos=tint_symbol_3};
   float const r = colors.col1;
   float const g = colors.col2;
   return;

test/bug/tint/744.wgsl.expected.msl

@@ -1,4 +1,4 @@
-SKIP: crbug.com/tint/817
+SKIP: crbug.com/tint/833 loop emission is broken
 
 #include <metal_stdlib>
 

test/bug/tint/824.wgsl.expected.msl

@@ -1,5 +1,3 @@
-SKIP: crbug.com/tint/817
-
 #include <metal_stdlib>
 
 using namespace metal;
@@ -8,23 +6,24 @@ struct Output {
   float4 color;
 };
 struct tint_symbol_2 {
-  uint VertexIndex [[vertex_id]];
-  uint InstanceIndex [[instance_id]];
-};
-struct tint_symbol_3 {
   float4 color [[user(locn0)]];
   float4 Position [[position]];
 };
+struct tint_array_wrapper_0 {
+  float2 array[4];
+};
+struct tint_array_wrapper_1 {
+  float4 array[4];
+};
 
-vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+vertex tint_symbol_2 tint_symbol(uint VertexIndex [[vertex_id]], uint InstanceIndex [[instance_id]]) {
-  uint const VertexIndex = tint_symbol_1.VertexIndex;
+  tint_array_wrapper_0 const zv = {float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)};
-  uint const InstanceIndex = tint_symbol_1.InstanceIndex;
+  float const z = zv.array[InstanceIndex].x;
-  float2 const zv[4] = {float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)};
-  float const z = zv[InstanceIndex].x;
   Output output = {};
   output.Position = float4(0.5f, 0.5f, z, 1.0f);
-  float4 const colors[4] = {float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)};
+  tint_array_wrapper_1 const colors = {float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)};
-  output.color = colors[InstanceIndex];
+  output.color = colors.array[InstanceIndex];
-  return {output.color, output.Position};
+  tint_symbol_2 const tint_symbol_3 = {.color=output.color, .Position=output.Position};
+  return tint_symbol_3;
 }
 

test/samples/compute_boids.wgsl.expected.msl

@@ -1,5 +1,3 @@
-SKIP: crbug.com/tint/817
-
 #include <metal_stdlib>
 
 using namespace metal;
@@ -27,11 +25,11 @@ struct SimParams {
   /* 0x0014 */ float rule2Scale;
   /* 0x0018 */ float rule3Scale;
 };
-struct Particles {
+struct tint_array_wrapper_0 {
-  /* 0x0000 */ Particle particles[5];
+  Particle array[5];
 };
-struct tint_symbol_5 {
+struct Particles {
-  uint3 gl_GlobalInvocationID [[thread_position_in_grid]];
+  /* 0x0000 */ tint_array_wrapper_0 particles;
 };
 
 vertex tint_symbol_2 vert_main(tint_symbol_1 tint_symbol [[stage_in]]) {
@@ -40,21 +38,22 @@ vertex tint_symbol_2 vert_main(tint_symbol_1 tint_symbol [[stage_in]]) {
   float2 const a_pos = tint_symbol.a_pos;
   float angle = -(  atan2(a_particleVel.x, a_particleVel.y));
   float2 pos = float2(((a_pos.x *   cos(angle)) - (a_pos.y *   sin(angle))), ((a_pos.x *   sin(angle)) + (a_pos.y *   cos(angle))));
-  return {float4((pos + a_particlePos), 0.0f, 1.0f)};
+  tint_symbol_2 const tint_symbol_5 = {.value=float4((pos + a_particlePos), 0.0f, 1.0f)};
+  return tint_symbol_5;
 }
 
 fragment tint_symbol_3 frag_main() {
-  return {float4(1.0f, 1.0f, 1.0f, 1.0f)};
+  tint_symbol_3 const tint_symbol_6 = {.value=float4(1.0f, 1.0f, 1.0f, 1.0f)};
+  return tint_symbol_6;
 }
 
-kernel void comp_main(tint_symbol_5 tint_symbol_4 [[stage_in]], constant SimParams& params [[buffer(0)]], device Particles& particlesA [[buffer(1)]], device Particles& particlesB [[buffer(2)]]) {
+kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], constant SimParams& params [[buffer(0)]], device Particles& particlesA [[buffer(1)]], device Particles& particlesB [[buffer(2)]]) {
-  uint3 const gl_GlobalInvocationID = tint_symbol_4.gl_GlobalInvocationID;
   uint index = gl_GlobalInvocationID.x;
   if ((index >= 5u)) {
     return;
   }
-  float2 vPos = particlesA.particles[index].pos;
+  float2 vPos = particlesA.particles.array[index].pos;
-  float2 vVel = particlesA.particles[index].vel;
+  float2 vVel = particlesA.particles.array[index].vel;
   float2 cMass = float2(0.0f, 0.0f);
   float2 cVel = float2(0.0f, 0.0f);
   float2 colVel = float2(0.0f, 0.0f);
@@ -78,8 +77,8 @@ kernel void comp_main(tint_symbol_5 tint_symbol_4 [[stage_in]], constant SimPara
         if ((i == index)) {
           continue;
         }
-        pos = particlesA.particles[i].pos.xy;
+        pos = particlesA.particles.array[i].pos.xy;
-        vel = particlesA.particles[i].vel.xy;
+        vel = particlesA.particles.array[i].vel.xy;
         if ((        distance(pos, vPos) < params.rule1Distance)) {
           cMass = (cMass + pos);
           cMassCount = (cMassCount + 1);
@@ -115,8 +114,8 @@ kernel void comp_main(tint_symbol_5 tint_symbol_4 [[stage_in]], constant SimPara
   if ((vPos.y > 1.0f)) {
     vPos.y = -1.0f;
   }
-  particlesB.particles[index].pos = vPos;
+  particlesB.particles.array[index].pos = vPos;
-  particlesB.particles[index].vel = vVel;
+  particlesB.particles.array[index].vel = vVel;
   return;
 }
 

test/samples/triangle.wgsl.expected.msl

@@ -1,25 +1,24 @@
-SKIP: crbug.com/tint/817
-
 #include <metal_stdlib>
 
 using namespace metal;
 struct tint_symbol_1 {
-  int VertexIndex [[vertex_id]];
-};
-struct tint_symbol_2 {
   float4 value [[position]];
 };
-struct tint_symbol_3 {
+struct tint_symbol_2 {
   float4 value [[color(0)]];
 };
+struct tint_array_wrapper_0 {
+  float2 array[3];
+};
 
-constant float2 pos[3] = {float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)};
+constant tint_array_wrapper_0 pos = {float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)};
-vertex tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol [[stage_in]]) {
+vertex tint_symbol_1 vtx_main(uint VertexIndex [[vertex_id]]) {
-  int const VertexIndex = tint_symbol.VertexIndex;
+  tint_symbol_1 const tint_symbol_3 = {.value=float4(pos.array[VertexIndex], 0.0f, 1.0f)};
-  return {float4(pos[VertexIndex], 0.0f, 1.0f)};
+  return tint_symbol_3;
 }
 
-fragment tint_symbol_3 frag_main() {
+fragment tint_symbol_2 frag_main() {
-  return {float4(1.0f, 0.0f, 0.0f, 1.0f)};
+  tint_symbol_2 const tint_symbol_4 = {.value=float4(1.0f, 0.0f, 0.0f, 1.0f)};
+  return tint_symbol_4;
 }
 

test/shader_io/compute_input_builtins.wgsl.expected.msl

@@ -1 +1,8 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(uint3 local_invocation_id [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], uint3 global_invocation_id [[thread_position_in_grid]], uint3 workgroup_id [[threadgroup_position_in_grid]]) {
+  uint const foo = (((local_invocation_id.x + local_invocation_index) + global_invocation_id.x) + workgroup_id.x);
+  return;
+}
+

test/shader_io/compute_input_builtins_struct.wgsl.expected.msl

@@ -1 +1,16 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+struct ComputeInputs {
+  uint3 local_invocation_id;
+  uint local_invocation_index;
+  uint3 global_invocation_id;
+  uint3 workgroup_id;
+};
+
+kernel void tint_symbol(uint3 tint_symbol_2 [[thread_position_in_threadgroup]], uint tint_symbol_3 [[thread_index_in_threadgroup]], uint3 tint_symbol_4 [[thread_position_in_grid]], uint3 tint_symbol_5 [[threadgroup_position_in_grid]]) {
+  ComputeInputs const inputs = {.local_invocation_id=tint_symbol_2, .local_invocation_index=tint_symbol_3, .global_invocation_id=tint_symbol_4, .workgroup_id=tint_symbol_5};
+  uint const foo = (((inputs.local_invocation_id.x + inputs.local_invocation_index) + inputs.global_invocation_id.x) + inputs.workgroup_id.x);
+  return;
+}
+

test/shader_io/compute_input_mixed.wgsl.expected.msl

@@ -1 +1,17 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+struct ComputeInputs0 {
+  uint3 local_invocation_id;
+};
+struct ComputeInputs1 {
+  uint3 workgroup_id;
+};
+
+kernel void tint_symbol(uint3 tint_symbol_2 [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], uint3 global_invocation_id [[thread_position_in_grid]], uint3 tint_symbol_3 [[threadgroup_position_in_grid]]) {
+  ComputeInputs0 const inputs0 = {.local_invocation_id=tint_symbol_2};
+  ComputeInputs1 const inputs1 = {.workgroup_id=tint_symbol_3};
+  uint const foo = (((inputs0.local_invocation_id.x + local_invocation_index) + global_invocation_id.x) + inputs1.workgroup_id.x);
+  return;
+}
+

test/shader_io/fragment_input_builtins.wgsl.expected.msl

@@ -1,18 +1,7 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_symbol_2 {
+fragment void tint_symbol(float4 position [[position]], bool front_facing [[front_facing]], uint sample_index [[sample_id]], uint sample_mask [[sample_mask]]) {
-  float4 position [[position]];
-  bool front_facing [[front_facing]];
-  uint sample_index [[sample_id]];
-  uint sample_mask [[sample_mask]];
-};
-
-fragment void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
-  float4 const position = tint_symbol_1.position;
-  bool const front_facing = tint_symbol_1.front_facing;
-  uint const sample_index = tint_symbol_1.sample_index;
-  uint const sample_mask = tint_symbol_1.sample_mask;
   if (front_facing) {
     float4 const foo = position;
     uint const bar = (sample_index + sample_mask);

test/shader_io/fragment_input_builtins_struct.wgsl.expected.msl

@@ -7,15 +7,9 @@ struct FragmentInputs {
   uint sample_index;
   uint sample_mask;
 };
-struct tint_symbol_2 {
-  float4 position [[position]];
-  bool front_facing [[front_facing]];
-  uint sample_index [[sample_id]];
-  uint sample_mask [[sample_mask]];
-};
 
-fragment void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+fragment void tint_symbol(float4 tint_symbol_2 [[position]], bool tint_symbol_3 [[front_facing]], uint tint_symbol_4 [[sample_id]], uint tint_symbol_5 [[sample_mask]]) {
-  FragmentInputs const inputs = {.position=tint_symbol_1.position, .front_facing=tint_symbol_1.front_facing, .sample_index=tint_symbol_1.sample_index, .sample_mask=tint_symbol_1.sample_mask};
+  FragmentInputs const inputs = {.position=tint_symbol_2, .front_facing=tint_symbol_3, .sample_index=tint_symbol_4, .sample_mask=tint_symbol_5};
   if (inputs.front_facing) {
     float4 const foo = inputs.position;
     uint const bar = (inputs.sample_index + inputs.sample_mask);

test/shader_io/fragment_input_mixed.wgsl.expected.msl

@@ -9,23 +9,17 @@ struct FragmentInputs1 {
   float4 loc3;
   uint sample_mask;
 };
-struct tint_symbol_2 {
+struct tint_symbol_4 {
   int loc0 [[user(locn0)]];
   uint loc1 [[user(locn1)]];
   float loc2 [[user(locn2)]];
   float4 loc3 [[user(locn3)]];
-  float4 position [[position]];
-  bool front_facing [[front_facing]];
-  uint sample_index [[sample_id]];
-  uint sample_mask [[sample_mask]];
 };
 
-fragment void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+fragment void tint_symbol(float4 tint_symbol_2 [[position]], bool front_facing [[front_facing]], uint sample_index [[sample_id]], uint tint_symbol_3 [[sample_mask]], tint_symbol_4 tint_symbol_1 [[stage_in]]) {
-  FragmentInputs0 const inputs0 = {.position=tint_symbol_1.position, .loc0=tint_symbol_1.loc0};
+  FragmentInputs0 const inputs0 = {.position=tint_symbol_2, .loc0=tint_symbol_1.loc0};
-  bool const front_facing = tint_symbol_1.front_facing;
   uint const loc1 = tint_symbol_1.loc1;
-  uint const sample_index = tint_symbol_1.sample_index;
+  FragmentInputs1 const inputs1 = {.loc3=tint_symbol_1.loc3, .sample_mask=tint_symbol_3};
-  FragmentInputs1 const inputs1 = {.loc3=tint_symbol_1.loc3, .sample_mask=tint_symbol_1.sample_mask};
   float const loc2 = tint_symbol_1.loc2;
   if (front_facing) {
     float4 const foo = inputs0.position;

test/shader_io/shared_struct_different_stages.wgsl.expected.msl

@@ -11,20 +11,19 @@ struct tint_symbol {
   float col2 [[user(locn2)]];
   float4 pos [[position]];
 };
-struct tint_symbol_3 {
+struct tint_symbol_4 {
   float col1 [[user(locn1)]];
   float col2 [[user(locn2)]];
-  float4 pos [[position]];
 };
 
 vertex tint_symbol vert_main() {
   Interface const tint_symbol_1 = {.col1=0.400000006f, .col2=0.600000024f, .pos=float4()};
-  tint_symbol const tint_symbol_4 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=tint_symbol_1.pos};
+  tint_symbol const tint_symbol_5 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=tint_symbol_1.pos};
-  return tint_symbol_4;
+  return tint_symbol_5;
 }
 
-fragment void frag_main(tint_symbol_3 tint_symbol_2 [[stage_in]]) {
+fragment void frag_main(float4 tint_symbol_3 [[position]], tint_symbol_4 tint_symbol_2 [[stage_in]]) {
-  Interface const colors = {.col1=tint_symbol_2.col1, .col2=tint_symbol_2.col2, .pos=tint_symbol_2.pos};
+  Interface const colors = {.col1=tint_symbol_2.col1, .col2=tint_symbol_2.col2, .pos=tint_symbol_3};
   float const r = colors.col1;
   float const g = colors.col2;
   return;

test/shader_io/shared_struct_storage_buffer.wgsl.expected.msl

@@ -8,14 +8,13 @@ struct S {
   /* 0x0080 */ packed_float4 v;
   /* 0x0090 */ int8_t tint_pad_1[112];
 };
-struct tint_symbol_1 {
+struct tint_symbol_2 {
   float f [[user(locn0)]];
   uint u [[user(locn1)]];
-  float4 v [[position]];
 };
 
-fragment void frag_main(tint_symbol_1 tint_symbol [[stage_in]], device S& output [[buffer(0)]]) {
+fragment void frag_main(float4 tint_symbol_1 [[position]], tint_symbol_2 tint_symbol [[stage_in]], device S& output [[buffer(0)]]) {
-  S const input = {.f=tint_symbol.f, .u=tint_symbol.u, .v=tint_symbol.v};
+  S const input = {.f=tint_symbol.f, .u=tint_symbol.u, .v=tint_symbol_1};
   float const f = input.f;
   uint const u = input.u;
   float4 const v = input.v;

test/shader_io/vertex_input_builtins.wgsl.expected.msl

@@ -1 +1,13 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_2 {
+  float4 value [[position]];
+};
+
+vertex tint_symbol_2 tint_symbol(uint vertex_index [[vertex_id]], uint instance_index [[instance_id]]) {
+  uint const foo = (vertex_index + instance_index);
+  tint_symbol_2 const tint_symbol_3 = {.value=float4()};
+  return tint_symbol_3;
+}
+

test/shader_io/vertex_input_builtins_struct.wgsl.expected.msl

@@ -1 +1,18 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+struct VertexInputs {
+  uint vertex_index;
+  uint instance_index;
+};
+struct tint_symbol_4 {
+  float4 value [[position]];
+};
+
+vertex tint_symbol_4 tint_symbol(uint tint_symbol_2 [[vertex_id]], uint tint_symbol_3 [[instance_id]]) {
+  VertexInputs const inputs = {.vertex_index=tint_symbol_2, .instance_index=tint_symbol_3};
+  uint const foo = (inputs.vertex_index + inputs.instance_index);
+  tint_symbol_4 const tint_symbol_5 = {.value=float4()};
+  return tint_symbol_5;
+}
+

test/shader_io/vertex_input_mixed.wgsl.expected.msl

@@ -1 +1,34 @@
-SKIP: crbug.com/tint/817 attribute only applies to parameters
+#include <metal_stdlib>
+
+using namespace metal;
+struct VertexInputs0 {
+  uint vertex_index;
+  int loc0;
+};
+struct VertexInputs1 {
+  float loc2;
+  float4 loc3;
+};
+struct tint_symbol_3 {
+  int loc0 [[attribute(0)]];
+  uint loc1 [[attribute(1)]];
+  float loc2 [[attribute(2)]];
+  float4 loc3 [[attribute(3)]];
+};
+struct tint_symbol_4 {
+  float4 value [[position]];
+};
+
+vertex tint_symbol_4 tint_symbol(uint tint_symbol_2 [[vertex_id]], uint instance_index [[instance_id]], tint_symbol_3 tint_symbol_1 [[stage_in]]) {
+  VertexInputs0 const inputs0 = {.vertex_index=tint_symbol_2, .loc0=tint_symbol_1.loc0};
+  uint const loc1 = tint_symbol_1.loc1;
+  VertexInputs1 const inputs1 = {.loc2=tint_symbol_1.loc2, .loc3=tint_symbol_1.loc3};
+  uint const foo = (inputs0.vertex_index + instance_index);
+  int const i = inputs0.loc0;
+  uint const u = loc1;
+  float const f = inputs1.loc2;
+  float4 const v = inputs1.loc3;
+  tint_symbol_4 const tint_symbol_5 = {.value=float4()};
+  return tint_symbol_5;
+}
+