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>
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 151ec79..763bb66 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -135,13 +135,14 @@
}
}
+ 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 @@
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 @@
}
// 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 @@
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;
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index d106b66..1bb9a32 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/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 @@
/// 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
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 8a81489..e089fca 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
@@ -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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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;
};
diff --git a/src/writer/hlsl/generator_impl_function_test.cc b/src/writer/hlsl/generator_impl_function_test.cc
index 5edfe2b..5f02e3a 100644
--- a/src/writer/hlsl/generator_impl_function_test.cc
+++ b/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 @@
)");
}
+// 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
diff --git a/src/writer/msl/generator_impl_function_test.cc b/src/writer/msl/generator_impl_function_test.cc
index 5cbe0d2..a5d871b 100644
--- a/src/writer/msl/generator_impl_function_test.cc
+++ b/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 @@
)");
}
+// 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
diff --git a/src/writer/spirv/builder_function_test.cc b/src/writer/spirv/builder_function_test.cc
index 266b7b7..ba4ad59 100644
--- a/src/writer/spirv/builder_function_test.cc
+++ b/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 @@
)");
}
+// 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
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index eee6edb..8266900 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/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::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 @@
} else if (type->IsVoid()) {
out_ << "void";
} else {
- error_ = "unknown type in EmitType";
+ error_ = "unknown type in EmitType: " + type->type_name();
return false;
}
diff --git a/src/writer/wgsl/generator_impl_function_test.cc b/src/writer/wgsl/generator_impl_function_test.cc
index 92f4b62..f7ba33b 100644
--- a/src/writer/wgsl/generator_impl_function_test.cc
+++ b/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 @@
)");
}
+// 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
