Fixup emitting duplicate globals in HLSL.

This CL fixes the issue with duplicate globals being emitted in HLSL if
used in multiple entry points. Tests are added for the other backends to
verify the issue does not exist there.

Bug: tint:297
Change-Id: I16d7504e8458fd375c6e1896758fe180ad963871
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/31880
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>

src/writer/hlsl/generator_impl.cc

@@ -135,13 +135,14 @@ bool GeneratorImpl::Generate(std::ostream& out) {
     }
   }
 
+  std::unordered_set<std::string> emitted_globals;
   // Make sure all entry point data is emitted before the entry point functions
   for (const auto& func : module_->functions()) {
     if (!func->IsEntryPoint()) {
       continue;
     }
 
-    if (!EmitEntryPointData(out, func.get())) {
+    if (!EmitEntryPointData(out, func.get(), emitted_globals)) {
       return false;
     }
   }
@@ -1136,7 +1137,10 @@ bool GeneratorImpl::EmitFunctionInternal(std::ostream& out,
   return true;
 }
 
-bool GeneratorImpl::EmitEntryPointData(std::ostream& out, ast::Function* func) {
+bool GeneratorImpl::EmitEntryPointData(
+    std::ostream& out,
+    ast::Function* func,
+    std::unordered_set<std::string>& emitted_globals) {
   std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> in_variables;
   std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> outvariables;
   for (auto data : func->referenced_location_variables()) {
@@ -1174,6 +1178,13 @@ bool GeneratorImpl::EmitEntryPointData(std::ostream& out, ast::Function* func) {
     }
     // auto* set = data.second.set;
 
+    // If the global has already been emitted we skip it, it's been emitted by
+    // a previous entry point.
+    if (emitted_globals.count(var->name()) != 0) {
+      continue;
+    }
+    emitted_globals.insert(var->name());
+
     auto* type = var->type()->UnwrapIfNeeded();
     if (type->IsStruct()) {
       auto* strct = type->AsStruct();
@@ -1210,6 +1221,13 @@ bool GeneratorImpl::EmitEntryPointData(std::ostream& out, ast::Function* func) {
     auto* var = data.first;
     auto* binding = data.second.binding;
 
+    // If the global has already been emitted we skip it, it's been emitted by
+    // a previous entry point.
+    if (emitted_globals.count(var->name()) != 0) {
+      continue;
+    }
+    emitted_globals.insert(var->name());
+
     if (!var->type()->IsAccessControl()) {
       error_ = "access control type required for storage buffer";
       return false;

src/writer/hlsl/generator_impl.h

@@ -17,6 +17,7 @@
 
 #include <string>
 #include <unordered_map>
+#include <unordered_set>
 
 #include "src/ast/intrinsic.h"
 #include "src/ast/literal.h"
@@ -193,8 +194,11 @@ class GeneratorImpl {
   /// Handles emitting information for an entry point
   /// @param out the output stream
   /// @param func the entry point
+  /// @param emitted_globals the set of globals emitted over all entry points
   /// @returns true if the entry point data was emitted
-  bool EmitEntryPointData(std::ostream& out, ast::Function* func);
+  bool EmitEntryPointData(std::ostream& out,
+                          ast::Function* func,
+                          std::unordered_set<std::string>& emitted_globals);
   /// Handles emitting the entry point function
   /// @param out the output stream
   /// @param func the entry point

src/writer/hlsl/generator_impl_function_entry_point_data_test.cc

@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include <memory>
+#include <unordered_set>
 
 #include "src/ast/assignment_statement.h"
 #include "src/ast/decorated_variable.h"
@@ -88,8 +89,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(result(), R"(struct vtx_main_in {
   float foo : TEXCOORD0;
   int bar : TEXCOORD1;
@@ -147,8 +151,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(result(), R"(struct vtx_main_out {
   float foo : TEXCOORD0;
   int bar : TEXCOORD1;
@@ -205,8 +212,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(result(), R"(struct main_in {
   float foo : TEXCOORD0;
   int bar : TEXCOORD1;
@@ -263,8 +273,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(result(), R"(struct main_out {
   float foo : SV_Target0;
   int bar : SV_Target1;
@@ -318,8 +331,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_FALSE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_FALSE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(gen().error(), R"(invalid location variable for pipeline stage)");
 }
 
@@ -368,8 +384,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_FALSE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_FALSE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(gen().error(), R"(invalid location variable for pipeline stage)");
 }
 
@@ -429,8 +448,11 @@ TEST_F(HlslGeneratorImplTest_EntryPoint,
 
   mod()->AddFunction(std::move(func));
 
+  std::unordered_set<std::string> globals;
+
   ASSERT_TRUE(td().Determine()) << td().error();
-  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr)) << gen().error();
+  ASSERT_TRUE(gen().EmitEntryPointData(out(), func_ptr, globals))
+      << gen().error();
   EXPECT_EQ(result(), R"(struct main_in {
   vector<float, 4> coord : SV_Position;
 };

src/writer/hlsl/generator_impl_function_test.cc

@@ -31,6 +31,7 @@
 #include "src/ast/sint_literal.h"
 #include "src/ast/stage_decoration.h"
 #include "src/ast/struct.h"
+#include "src/ast/struct_block_decoration.h"
 #include "src/ast/struct_member_offset_decoration.h"
 #include "src/ast/type/access_control_type.h"
 #include "src/ast/type/array_type.h"
@@ -1133,6 +1134,121 @@ TEST_F(HlslGeneratorImplTest_Function, Emit_Function_WithArrayParams) {
 )");
 }
 
+// https://crbug.com/tint/297
+TEST_F(HlslGeneratorImplTest_Function,
+       Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
+  // [[block]] struct Data {
+  //   [[offset(0)]] d : f32;
+  // };
+  // [[binding(0), set(0)]] var<storage_buffer> data : Data;
+  //
+  // [[stage(compute)]]
+  // fn a() -> void {
+  //   return;
+  // }
+  //
+  // [[stage(compute)]]
+  // fn b() -> void {
+  //   return;
+  // }
+
+  ast::type::VoidType void_type;
+  ast::type::F32Type f32;
+
+  ast::StructMemberList members;
+  ast::StructMemberDecorationList a_deco;
+  a_deco.push_back(
+      std::make_unique<ast::StructMemberOffsetDecoration>(0, Source{}));
+  members.push_back(
+      std::make_unique<ast::StructMember>("d", &f32, std::move(a_deco)));
+
+  ast::StructDecorationList s_decos;
+  s_decos.push_back(std::make_unique<ast::StructBlockDecoration>(Source{}));
+
+  auto str =
+      std::make_unique<ast::Struct>(std::move(s_decos), std::move(members));
+
+  ast::type::StructType s("Data", std::move(str));
+  ast::type::AccessControlType ac(ast::AccessControl::kReadWrite, &s);
+
+  auto data_var =
+      std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
+          "data", ast::StorageClass::kStorageBuffer, &ac));
+
+  ast::VariableDecorationList decos;
+  decos.push_back(std::make_unique<ast::BindingDecoration>(0, Source{}));
+  decos.push_back(std::make_unique<ast::SetDecoration>(0, Source{}));
+  data_var->set_decorations(std::move(decos));
+
+  mod()->AddConstructedType(&s);
+  td().RegisterVariableForTesting(data_var.get());
+  mod()->AddGlobalVariable(std::move(data_var));
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("a", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod()->AddFunction(std::move(func));
+  }
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("b", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod()->AddFunction(std::move(func));
+  }
+
+  ASSERT_TRUE(td().Determine()) << td().error();
+  ASSERT_TRUE(gen().Generate(out())) << gen().error();
+  EXPECT_EQ(result(), R"(struct Data {
+  float d;
+};
+
+RWByteAddressBuffer data : register(u0);
+
+[numthreads(1, 1, 1)]
+void a() {
+  float v = asfloat(data.Load(0));
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void b() {
+  float v = asfloat(data.Load(0));
+  return;
+}
+
+)");
+}
+
 }  // namespace
 }  // namespace hlsl
 }  // namespace writer

src/writer/msl/generator_impl_function_test.cc

@@ -32,6 +32,7 @@
 #include "src/ast/sint_literal.h"
 #include "src/ast/stage_decoration.h"
 #include "src/ast/struct.h"
+#include "src/ast/struct_block_decoration.h"
 #include "src/ast/struct_member.h"
 #include "src/ast/struct_member_decoration.h"
 #include "src/ast/struct_member_offset_decoration.h"
@@ -1168,6 +1169,126 @@ TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) {
 )");
 }
 
+// https://crbug.com/tint/297
+TEST_F(MslGeneratorImplTest,
+       Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar) {
+  // [[block]] struct Data {
+  //   [[offset(0)]] d : f32;
+  // };
+  // [[binding(0), set(0)]] var<storage_buffer> data : Data;
+  //
+  // [[stage(compute)]]
+  // fn a() -> void {
+  //   return;
+  // }
+  //
+  // [[stage(compute)]]
+  // fn b() -> void {
+  //   return;
+  // }
+
+  ast::type::VoidType void_type;
+  ast::type::F32Type f32;
+
+  ast::StructMemberList members;
+  ast::StructMemberDecorationList a_deco;
+  a_deco.push_back(
+      std::make_unique<ast::StructMemberOffsetDecoration>(0, Source{}));
+  members.push_back(
+      std::make_unique<ast::StructMember>("d", &f32, std::move(a_deco)));
+
+  ast::StructDecorationList s_decos;
+  s_decos.push_back(std::make_unique<ast::StructBlockDecoration>(Source{}));
+
+  auto str =
+      std::make_unique<ast::Struct>(std::move(s_decos), std::move(members));
+
+  ast::type::StructType s("Data", std::move(str));
+  ast::type::AccessControlType ac(ast::AccessControl::kReadWrite, &s);
+
+  auto data_var =
+      std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
+          "data", ast::StorageClass::kStorageBuffer, &ac));
+
+  ast::VariableDecorationList decos;
+  decos.push_back(std::make_unique<ast::BindingDecoration>(0, Source{}));
+  decos.push_back(std::make_unique<ast::SetDecoration>(0, Source{}));
+  data_var->set_decorations(std::move(decos));
+
+  Context ctx;
+  ast::Module mod;
+  TypeDeterminer td(&ctx, &mod);
+
+  mod.AddConstructedType(&s);
+
+  td.RegisterVariableForTesting(data_var.get());
+  mod.AddGlobalVariable(std::move(data_var));
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("a", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("b", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  ASSERT_TRUE(td.Determine()) << td.error();
+
+  GeneratorImpl g(&mod);
+  ASSERT_TRUE(g.Generate()) << g.error();
+  EXPECT_EQ(g.result(), R"(#include <metal_stdlib>
+
+struct Data {
+  float d;
+};
+
+kernel void a(device Data& data [[buffer(0)]]) {
+  float v = data.d;
+  return;
+}
+
+kernel void b(device Data& data [[buffer(0)]]) {
+  float v = data.d;
+  return;
+}
+
+)");
+}
+
 }  // namespace
 }  // namespace msl
 }  // namespace writer

src/writer/spirv/builder_function_test.cc

@@ -17,13 +17,22 @@
 #include "gtest/gtest.h"
 #include "spirv/unified1/spirv.h"
 #include "spirv/unified1/spirv.hpp11"
+#include "src/ast/decorated_variable.h"
 #include "src/ast/function.h"
 #include "src/ast/identifier_expression.h"
+#include "src/ast/member_accessor_expression.h"
 #include "src/ast/return_statement.h"
+#include "src/ast/stage_decoration.h"
+#include "src/ast/struct.h"
+#include "src/ast/struct_block_decoration.h"
+#include "src/ast/struct_member_offset_decoration.h"
+#include "src/ast/type/access_control_type.h"
 #include "src/ast/type/f32_type.h"
 #include "src/ast/type/i32_type.h"
+#include "src/ast/type/struct_type.h"
 #include "src/ast/type/void_type.h"
 #include "src/ast/variable.h"
+#include "src/ast/variable_decl_statement.h"
 #include "src/context.h"
 #include "src/type_determiner.h"
 #include "src/writer/spirv/builder.h"
@@ -150,6 +159,155 @@ TEST_F(BuilderTest, FunctionType_DeDuplicate) {
 )");
 }
 
+// https://crbug.com/tint/297
+TEST_F(BuilderTest, Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
+  // [[block]] struct Data {
+  //   [[offset(0)]] d : f32;
+  // };
+  // [[binding(0), set(0)]] var<storage_buffer> data : Data;
+  //
+  // [[stage(compute)]]
+  // fn a() -> void {
+  //   return;
+  // }
+  //
+  // [[stage(compute)]]
+  // fn b() -> void {
+  //   return;
+  // }
+
+  ast::type::VoidType void_type;
+  ast::type::F32Type f32;
+
+  ast::StructMemberList members;
+  ast::StructMemberDecorationList a_deco;
+  a_deco.push_back(
+      std::make_unique<ast::StructMemberOffsetDecoration>(0, Source{}));
+  members.push_back(
+      std::make_unique<ast::StructMember>("d", &f32, std::move(a_deco)));
+
+  ast::StructDecorationList s_decos;
+  s_decos.push_back(std::make_unique<ast::StructBlockDecoration>(Source{}));
+
+  auto str =
+      std::make_unique<ast::Struct>(std::move(s_decos), std::move(members));
+
+  ast::type::StructType s("Data", std::move(str));
+  ast::type::AccessControlType ac(ast::AccessControl::kReadWrite, &s);
+
+  auto data_var =
+      std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
+          "data", ast::StorageClass::kStorageBuffer, &ac));
+
+  ast::VariableDecorationList decos;
+  decos.push_back(std::make_unique<ast::BindingDecoration>(0, Source{}));
+  decos.push_back(std::make_unique<ast::SetDecoration>(0, Source{}));
+  data_var->set_decorations(std::move(decos));
+
+  Context ctx;
+  ast::Module mod;
+  TypeDeterminer td(&ctx, &mod);
+
+  mod.AddConstructedType(&s);
+
+  td.RegisterVariableForTesting(data_var.get());
+  mod.AddGlobalVariable(std::move(data_var));
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("a", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("b", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  ASSERT_TRUE(td.Determine()) << td.error();
+
+  Builder b(&mod);
+  ASSERT_TRUE(b.Build());
+  EXPECT_EQ(DumpBuilder(b), R"(OpCapability Shader
+OpCapability VulkanMemoryModel
+OpExtension "SPV_KHR_vulkan_memory_model"
+OpMemoryModel Logical Vulkan
+OpEntryPoint GLCompute %7 "a"
+OpEntryPoint GLCompute %17 "b"
+OpExecutionMode %7 LocalSize 1 1 1
+OpExecutionMode %17 LocalSize 1 1 1
+OpName %3 "Data"
+OpMemberName %3 0 "d"
+OpName %1 "data"
+OpName %7 "a"
+OpName %13 "v"
+OpName %17 "b"
+OpName %20 "v"
+OpDecorate %3 Block
+OpMemberDecorate %3 0 Offset 0
+OpDecorate %1 Binding 0
+OpDecorate %1 DescriptorSet 0
+%4 = OpTypeFloat 32
+%3 = OpTypeStruct %4
+%2 = OpTypePointer StorageBuffer %3
+%1 = OpVariable %2 StorageBuffer
+%6 = OpTypeVoid
+%5 = OpTypeFunction %6
+%9 = OpTypeInt 32 0
+%10 = OpConstant %9 0
+%11 = OpTypePointer StorageBuffer %4
+%14 = OpTypePointer Function %4
+%15 = OpConstantNull %4
+%7 = OpFunction %6 None %5
+%8 = OpLabel
+%13 = OpVariable %14 Function %15
+%12 = OpAccessChain %11 %1 %10
+%16 = OpLoad %4 %12
+OpStore %13 %16
+OpReturn
+OpFunctionEnd
+%17 = OpFunction %6 None %5
+%18 = OpLabel
+%20 = OpVariable %14 Function %15
+%19 = OpAccessChain %11 %1 %10
+%21 = OpLoad %4 %19
+OpStore %20 %21
+OpReturn
+OpFunctionEnd
+)");
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer

src/writer/wgsl/generator_impl.cc

@@ -51,6 +51,7 @@
 #include "src/ast/struct_member.h"
 #include "src/ast/struct_member_offset_decoration.h"
 #include "src/ast/switch_statement.h"
+#include "src/ast/type/access_control_type.h"
 #include "src/ast/type/array_type.h"
 #include "src/ast/type/depth_texture_type.h"
 #include "src/ast/type/matrix_type.h"
@@ -401,7 +402,24 @@ bool GeneratorImpl::EmitImageFormat(const ast::type::ImageFormat fmt) {
 }
 
 bool GeneratorImpl::EmitType(ast::type::Type* type) {
-  if (type->IsAlias()) {
+  if (type->IsAccessControl()) {
+    auto* ac = type->AsAccessControl();
+    // TODO(dsinclair): Access control isn't supported in WGSL yet, so this
+    // is disabled for now.
+    //
+    // out_ << "[[access(";
+    // if (ac->IsReadOnly()) {
+    //   out_ << "read";
+    // } else if (ac->IsWriteOnly()) {
+    //   out_ << "write";
+    // } else {
+    //   out_ << "read_write";
+    // }
+    // out_ << ")]]" << std::endl;
+    if (!EmitType(ac->type())) {
+      return false;
+    }
+  } else if (type->IsAlias()) {
     out_ << type->AsAlias()->name();
   } else if (type->IsArray()) {
     auto* ary = type->AsArray();
@@ -544,7 +562,7 @@ bool GeneratorImpl::EmitType(ast::type::Type* type) {
   } else if (type->IsVoid()) {
     out_ << "void";
   } else {
-    error_ = "unknown type in EmitType";
+    error_ = "unknown type in EmitType: " + type->type_name();
     return false;
   }
 

src/writer/wgsl/generator_impl_function_test.cc

@@ -13,16 +13,25 @@
 // limitations under the License.
 
 #include "gtest/gtest.h"
+#include "src/ast/decorated_variable.h"
 #include "src/ast/discard_statement.h"
 #include "src/ast/function.h"
+#include "src/ast/member_accessor_expression.h"
+#include "src/ast/module.h"
 #include "src/ast/pipeline_stage.h"
 #include "src/ast/return_statement.h"
 #include "src/ast/stage_decoration.h"
+#include "src/ast/struct_block_decoration.h"
+#include "src/ast/struct_member_offset_decoration.h"
+#include "src/ast/type/access_control_type.h"
 #include "src/ast/type/f32_type.h"
 #include "src/ast/type/i32_type.h"
 #include "src/ast/type/void_type.h"
 #include "src/ast/variable.h"
+#include "src/ast/variable_decl_statement.h"
 #include "src/ast/workgroup_decoration.h"
+#include "src/context.h"
+#include "src/type_determiner.h"
 #include "src/writer/wgsl/generator_impl.h"
 
 namespace tint {
@@ -152,6 +161,130 @@ TEST_F(WgslGeneratorImplTest, Emit_Function_WithDecoration_Multiple) {
 )");
 }
 
+// https://crbug.com/tint/297
+TEST_F(WgslGeneratorImplTest,
+       Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar) {
+  // [[block]] struct Data {
+  //   [[offset(0)]] d : f32;
+  // };
+  // [[binding(0), set(0)]] var<storage_buffer> data : Data;
+  //
+  // [[stage(compute)]]
+  // fn a() -> void {
+  //   return;
+  // }
+  //
+  // [[stage(compute)]]
+  // fn b() -> void {
+  //   return;
+  // }
+
+  ast::type::VoidType void_type;
+  ast::type::F32Type f32;
+
+  ast::StructMemberList members;
+  ast::StructMemberDecorationList a_deco;
+  a_deco.push_back(
+      std::make_unique<ast::StructMemberOffsetDecoration>(0, Source{}));
+  members.push_back(
+      std::make_unique<ast::StructMember>("d", &f32, std::move(a_deco)));
+
+  ast::StructDecorationList s_decos;
+  s_decos.push_back(std::make_unique<ast::StructBlockDecoration>(Source{}));
+
+  auto str =
+      std::make_unique<ast::Struct>(std::move(s_decos), std::move(members));
+
+  ast::type::StructType s("Data", std::move(str));
+  ast::type::AccessControlType ac(ast::AccessControl::kReadWrite, &s);
+
+  auto data_var =
+      std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
+          "data", ast::StorageClass::kStorageBuffer, &ac));
+
+  ast::VariableDecorationList decos;
+  decos.push_back(std::make_unique<ast::BindingDecoration>(0, Source{}));
+  decos.push_back(std::make_unique<ast::SetDecoration>(0, Source{}));
+  data_var->set_decorations(std::move(decos));
+
+  Context ctx;
+  ast::Module mod;
+  TypeDeterminer td(&ctx, &mod);
+
+  mod.AddConstructedType(&s);
+
+  td.RegisterVariableForTesting(data_var.get());
+  mod.AddGlobalVariable(std::move(data_var));
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("a", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  {
+    ast::VariableList params;
+    auto func =
+        std::make_unique<ast::Function>("b", std::move(params), &void_type);
+    func->add_decoration(std::make_unique<ast::StageDecoration>(
+        ast::PipelineStage::kCompute, Source{}));
+
+    auto var = std::make_unique<ast::Variable>(
+        "v", ast::StorageClass::kFunction, &f32);
+    var->set_constructor(std::make_unique<ast::MemberAccessorExpression>(
+        std::make_unique<ast::IdentifierExpression>("data"),
+        std::make_unique<ast::IdentifierExpression>("d")));
+
+    auto body = std::make_unique<ast::BlockStatement>();
+    body->append(std::make_unique<ast::VariableDeclStatement>(std::move(var)));
+    body->append(std::make_unique<ast::ReturnStatement>());
+    func->set_body(std::move(body));
+
+    mod.AddFunction(std::move(func));
+  }
+
+  ASSERT_TRUE(td.Determine()) << td.error();
+
+  GeneratorImpl g;
+  ASSERT_TRUE(g.Generate(mod)) << g.error();
+  EXPECT_EQ(g.result(), R"([[block]]
+struct Data {
+  [[offset(0)]]
+  d : f32;
+};
+
+[[binding(0), set(0)]] var<storage_buffer> data : Data;
+
+[[stage(compute)]]
+fn a() -> void {
+  var v : f32 = data.d;
+  return;
+}
+
+[[stage(compute)]]
+fn b() -> void {
+  var v : f32 = data.d;
+  return;
+}
+
+)");
+}
+
 }  // namespace
 }  // namespace wgsl
 }  // namespace writer