encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-05-13 10:51:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove EntryPoint.

Browse Source 
This CL removes the EntryPoint node and transitions everything to the
stage decoration.

Change-Id: Ib2840155905c8fa60ff35870f0c4b6705efb73ff
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/28705
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: David Neto <dneto@google.com>
Reviewed-by: Sarah Mashayekhi <sarahmashay@google.com>
This commit is contained in:
dan sinclair 2020-09-22 14:53:03 +00:00 committed by Commit Bot service account
parent de2a019a7f
commit 5f8126271d
55 changed files with 86 additions and 4856 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
BUILD.gn
View File

@ -255,8 +255,6 @@ source_set("libtint_core_src") {
"src/ast/discard_statement.h",
"src/ast/else_statement.cc",
"src/ast/else_statement.h",
"src/ast/entry_point.cc",
"src/ast/entry_point.h",
"src/ast/expression.cc",
"src/ast/expression.h",
"src/ast/fallthrough_statement.cc",
@ -706,7 +704,6 @@ source_set("tint_unittests_core_src") {
"src/ast/decorated_variable_test.cc",
"src/ast/discard_statement_test.cc",
"src/ast/else_statement_test.cc",
"src/ast/entry_point_test.cc",
"src/ast/expression_test.cc",
"src/ast/fallthrough_statement_test.cc",
"src/ast/float_literal_test.cc",
@ -840,7 +837,6 @@ source_set("tint_unittests_spv_writer_src") {
"src/writer/spirv/builder_cast_expression_test.cc",
"src/writer/spirv/builder_constructor_expression_test.cc",
"src/writer/spirv/builder_discard_test.cc",
"src/writer/spirv/builder_entry_point_test.cc",
"src/writer/spirv/builder_format_conversion_test.cc",
"src/writer/spirv/builder_function_decoration_test.cc",
"src/writer/spirv/builder_function_test.cc",
@ -899,7 +895,6 @@ source_set("tint_unittests_wgsl_reader_src") {
"src/reader/wgsl/parser_impl_depth_texture_type_test.cc",
"src/reader/wgsl/parser_impl_else_stmt_test.cc",
"src/reader/wgsl/parser_impl_elseif_stmt_test.cc",
"src/reader/wgsl/parser_impl_entry_point_decl_test.cc",
"src/reader/wgsl/parser_impl_equality_expression_test.cc",
"src/reader/wgsl/parser_impl_exclusive_or_expression_test.cc",
"src/reader/wgsl/parser_impl_for_stmt_test.cc",
@ -992,7 +987,6 @@ source_set("tint_unittests_wgsl_writer_src") {
"src/writer/wgsl/generator_impl_constructor_test.cc",
"src/writer/wgsl/generator_impl_continue_test.cc",
"src/writer/wgsl/generator_impl_discard_test.cc",
"src/writer/wgsl/generator_impl_entry_point_test.cc",
"src/writer/wgsl/generator_impl_fallthrough_test.cc",
"src/writer/wgsl/generator_impl_function_test.cc",
"src/writer/wgsl/generator_impl_identifier_test.cc",
@ -1043,7 +1037,6 @@ source_set("tint_unittests_msl_writer_src") {
"src/writer/msl/generator_impl_constructor_test.cc",
"src/writer/msl/generator_impl_continue_test.cc",
"src/writer/msl/generator_impl_discard_test.cc",
"src/writer/msl/generator_impl_entry_point_test.cc",
"src/writer/msl/generator_impl_function_entry_point_data_test.cc",
"src/writer/msl/generator_impl_function_test.cc",
"src/writer/msl/generator_impl_identifier_test.cc",
@ -1096,7 +1089,6 @@ source_set("tint_unittests_hlsl_writer_src") {
"src/writer/hlsl/generator_impl_constructor_test.cc",
"src/writer/hlsl/generator_impl_continue_test.cc",
"src/writer/hlsl/generator_impl_discard_test.cc",
"src/writer/hlsl/generator_impl_entry_point_test.cc",
"src/writer/hlsl/generator_impl_function_entry_point_data_test.cc",
"src/writer/hlsl/generator_impl_function_test.cc",
"src/writer/hlsl/generator_impl_identifier_test.cc",

8
src/CMakeLists.txt
View File

@ -76,8 +76,6 @@ set(TINT_LIB_SRCS
ast/discard_statement.h
ast/else_statement.cc
ast/else_statement.h
ast/entry_point.cc
ast/entry_point.h
ast/expression.cc
ast/expression.h
ast/fallthrough_statement.cc
@ -315,7 +313,6 @@ set(TINT_TEST_SRCS
ast/discard_statement_test.cc
ast/decorated_variable_test.cc
ast/else_statement_test.cc
ast/entry_point_test.cc
ast/expression_test.cc
ast/fallthrough_statement_test.cc
ast/float_literal_test.cc
@ -429,7 +426,6 @@ if(${TINT_BUILD_WGSL_READER})
reader/wgsl/parser_impl_depth_texture_type_test.cc
reader/wgsl/parser_impl_else_stmt_test.cc
reader/wgsl/parser_impl_elseif_stmt_test.cc
reader/wgsl/parser_impl_entry_point_decl_test.cc
reader/wgsl/parser_impl_equality_expression_test.cc
reader/wgsl/parser_impl_exclusive_or_expression_test.cc
reader/wgsl/parser_impl_for_stmt_test.cc
@ -500,7 +496,6 @@ if(${TINT_BUILD_SPV_WRITER})
writer/spirv/builder_cast_expression_test.cc
writer/spirv/builder_constructor_expression_test.cc
writer/spirv/builder_discard_test.cc
writer/spirv/builder_entry_point_test.cc
writer/spirv/builder_format_conversion_test.cc
writer/spirv/builder_function_decoration_test.cc
writer/spirv/builder_function_test.cc
@ -539,7 +534,6 @@ if(${TINT_BUILD_WGSL_WRITER})
writer/wgsl/generator_impl_constructor_test.cc
writer/wgsl/generator_impl_continue_test.cc
writer/wgsl/generator_impl_discard_test.cc
writer/wgsl/generator_impl_entry_point_test.cc
writer/wgsl/generator_impl_fallthrough_test.cc
writer/wgsl/generator_impl_function_test.cc
writer/wgsl/generator_impl_identifier_test.cc
@ -571,7 +565,6 @@ if(${TINT_BUILD_MSL_WRITER})
writer/msl/generator_impl_constructor_test.cc
writer/msl/generator_impl_continue_test.cc
writer/msl/generator_impl_discard_test.cc
writer/msl/generator_impl_entry_point_test.cc
writer/msl/generator_impl_function_entry_point_data_test.cc
writer/msl/generator_impl_function_test.cc
writer/msl/generator_impl_identifier_test.cc
@ -606,7 +599,6 @@ if (${TINT_BUILD_HLSL_WRITER})
writer/hlsl/generator_impl_constructor_test.cc
writer/hlsl/generator_impl_continue_test.cc
writer/hlsl/generator_impl_discard_test.cc
writer/hlsl/generator_impl_entry_point_test.cc
writer/hlsl/generator_impl_function_entry_point_data_test.cc
writer/hlsl/generator_impl_function_test.cc
writer/hlsl/generator_impl_identifier_test.cc

53
src/ast/entry_point.cc
View File

@ -1,53 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/ast/entry_point.h"
namespace tint {
namespace ast {
EntryPoint::EntryPoint(PipelineStage stage,
const std::string& name,
const std::string& fn_name)
: Node(), stage_(stage), name_(name), fn_name_(fn_name) {}
EntryPoint::EntryPoint(const Source& source,
PipelineStage stage,
const std::string& name,
const std::string& fn_name)
: Node(source), stage_(stage), name_(name), fn_name_(fn_name) {}
EntryPoint::~EntryPoint() = default;
bool EntryPoint::IsValid() const {
if (stage_ == PipelineStage::kNone) {
return false;
}
if (fn_name_.length() == 0) {
return false;
}
return true;
}
void EntryPoint::to_str(std::ostream& out, size_t indent) const {
make_indent(out, indent);
out << "EntryPoint{" << stage_;
if (name_.length() > 0)
out << " as " << name_;
out << " = " << fn_name_ << "}" << std::endl;
}
} // namespace ast
} // namespace tint

94
src/ast/entry_point.h
View File

@ -1,94 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef SRC_AST_ENTRY_POINT_H_
#define SRC_AST_ENTRY_POINT_H_
#include <memory>
#include <ostream>
#include <string>
#include <vector>
#include "src/ast/node.h"
#include "src/ast/pipeline_stage.h"
namespace tint {
namespace ast {
/// An entry point statement.
class EntryPoint : public Node {
public:
/// Constructor
EntryPoint() = default;
/// Constructor
/// @param stage the entry point stage
/// @param name the entry point name
/// @param fn_name the function name
EntryPoint(PipelineStage stage,
const std::string& name,
const std::string& fn_name);
/// Constructor
/// @param source the source of the entry point
/// @param stage the entry point stage
/// @param name the entry point name
/// @param fn_name the function name
EntryPoint(const Source& source,
PipelineStage stage,
const std::string& name,
const std::string& fn_name);
/// Move constructor
EntryPoint(EntryPoint&&) = default;
~EntryPoint() override;
/// Sets the entry point name
/// @param name the name to set
void set_name(const std::string& name) { name_ = name; }
/// @returns the entry points name
const std::string& name() const { return name_; }
/// Sets the entry point function name
/// @param name the function name
void set_function_name(const std::string& name) { fn_name_ = name; }
/// @returns the function name for the entry point
const std::string& function_name() const { return fn_name_; }
/// Sets the piepline stage
/// @param stage the stage to set
void set_pipeline_stage(PipelineStage stage) { stage_ = stage; }
/// @returns the pipeline stage for the entry point
PipelineStage stage() const { return stage_; }
/// @returns true if the entry point is valid
bool IsValid() const override;
/// Writes a representation of the entry point to the output stream
/// @param out the stream to write too
/// @param indent number of spaces to ident the node when writing
void to_str(std::ostream& out, size_t indent) const override;
private:
EntryPoint(const EntryPoint&) = delete;
Source source_;
PipelineStage stage_;
std::string name_;
std::string fn_name_;
};
/// A list of unique entry points.
using EntryPointList = std::vector<std::unique_ptr<EntryPoint>>;
} // namespace ast
} // namespace tint
#endif // SRC_AST_ENTRY_POINT_H_

101
src/ast/entry_point_test.cc
View File

@ -1,101 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/ast/entry_point.h"
#include <sstream>
#include "gtest/gtest.h"
namespace tint {
namespace ast {
namespace {
using EntryPointTest = testing::Test;
TEST_F(EntryPointTest, Creation) {
EntryPoint e(PipelineStage::kVertex, "main", "vtx_main");
EXPECT_EQ(e.name(), "main");
EXPECT_EQ(e.function_name(), "vtx_main");
EXPECT_EQ(e.stage(), PipelineStage::kVertex);
EXPECT_EQ(e.line(), 0u);
EXPECT_EQ(e.column(), 0u);
}
TEST_F(EntryPointTest, CreationWithSource) {
Source s{27, 4};
EntryPoint e(s, PipelineStage::kVertex, "main", "vtx_main");
EXPECT_EQ(e.name(), "main");
EXPECT_EQ(e.function_name(), "vtx_main");
EXPECT_EQ(e.stage(), PipelineStage::kVertex);
EXPECT_EQ(e.line(), 27u);
EXPECT_EQ(e.column(), 4u);
}
TEST_F(EntryPointTest, CreationEmpty) {
Source s{27, 4};
EntryPoint e;
e.set_source(s);
e.set_pipeline_stage(PipelineStage::kFragment);
e.set_function_name("my_func");
e.set_name("a_name");
EXPECT_EQ(e.function_name(), "my_func");
EXPECT_EQ(e.name(), "a_name");
EXPECT_EQ(e.stage(), PipelineStage::kFragment);
EXPECT_EQ(e.line(), 27u);
EXPECT_EQ(e.column(), 4u);
}
TEST_F(EntryPointTest, IsValid) {
EntryPoint e(PipelineStage::kVertex, "", "vtx_main");
EXPECT_TRUE(e.IsValid());
}
TEST_F(EntryPointTest, IsValid_MissingFunctionName) {
EntryPoint e(PipelineStage::kVertex, "main", "");
EXPECT_FALSE(e.IsValid());
}
TEST_F(EntryPointTest, IsValid_MissingStage) {
EntryPoint e(PipelineStage::kNone, "main", "fn");
EXPECT_FALSE(e.IsValid());
}
TEST_F(EntryPointTest, IsValid_MissingBoth) {
EntryPoint e;
EXPECT_FALSE(e.IsValid());
}
TEST_F(EntryPointTest, ToStr) {
EntryPoint e(PipelineStage::kVertex, "text", "vtx_main");
std::ostringstream out;
e.to_str(out, 2);
EXPECT_EQ(out.str(), R"( EntryPoint{vertex as text = vtx_main}
)");
}
TEST_F(EntryPointTest, ToStr_NoName) {
EntryPoint e(PipelineStage::kVertex, "", "vtx_main");
std::ostringstream out;
e.to_str(out, 2);
EXPECT_EQ(out.str(), R"( EntryPoint{vertex = vtx_main}
)");
}
} // namespace
} // namespace ast
} // namespace tint

20
src/ast/module.cc
View File

@ -45,12 +45,14 @@ Function* Module::FindFunctionByName(const std::string& name) const {
return nullptr;
}
bool Module::IsFunctionEntryPoint(const std::string& name) const {
for (const auto& ep : entry_points_) {
if (ep->function_name() == name)
return true;
Function* Module::FindFunctionByNameAndStage(const std::string& name,
ast::PipelineStage stage) const {
for (const auto& func : functions_) {
if (func->name() == name && func->pipeline_stage() == stage) {
return func.get();
}
}
return false;
return nullptr;
}
bool Module::IsValid() const {
@ -64,11 +66,6 @@ bool Module::IsValid() const {
return false;
}
}
for (const auto& ep : entry_points_) {
if (ep == nullptr || !ep->IsValid()) {
return false;
}
}
for (auto* const alias : alias_types_) {
if (alias == nullptr) {
return false;
@ -93,9 +90,6 @@ std::string Module::to_str() const {
for (const auto& var : global_variables_) {
var->to_str(out, indent);
}
for (const auto& ep : entry_points_) {
ep->to_str(out, indent);
}
for (auto* const alias : alias_types_) {
for (size_t i = 0; i < indent; ++i) {
out << " ";

21
src/ast/module.h
View File

@ -20,7 +20,6 @@
#include <utility>
#include <vector>
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/type/alias_type.h"
@ -60,19 +59,6 @@ class Module {
/// @returns the global variables for the module
VariableList& global_variables() { return global_variables_; }
/// Adds an entry point to the module
/// @param ep the entry point to add
void AddEntryPoint(std::unique_ptr<EntryPoint> ep) {
entry_points_.push_back(std::move(ep));
}
/// @returns the entry points in the module
const EntryPointList& entry_points() const { return entry_points_; }
/// Checks if the given function name is an entry point function
/// @param name the function name
/// @returns true if name is an entry point function
bool IsFunctionEntryPoint(const std::string& name) const;
/// Adds a type alias to the module
/// @param type the alias to add
void AddAliasType(type::AliasType* type) { alias_types_.push_back(type); }
@ -92,6 +78,12 @@ class Module {
/// @param name the name to search for
/// @returns the associated function or nullptr if none exists
Function* FindFunctionByName(const std::string& name) const;
/// Returns the function with the given name
/// @param name the name to search for
/// @param stage the pipeline stage
/// @returns the associated function or nullptr if none exists
Function* FindFunctionByNameAndStage(const std::string& name,
ast::PipelineStage stage) const;
/// @returns true if all required fields in the AST are present.
bool IsValid() const;
@ -104,7 +96,6 @@ class Module {
ImportList imports_;
VariableList global_variables_;
EntryPointList entry_points_;
// The alias types are owned by the type manager
std::vector<type::AliasType*> alias_types_;
FunctionList functions_;

33
src/ast/module_test.cc
View File

@ -18,7 +18,6 @@
#include <utility>
#include "gmock/gmock.h"
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/type/f32_type.h"
@ -91,19 +90,6 @@ TEST_F(ModuleTest, LookupFunction) {
EXPECT_EQ(func_ptr, m.FindFunctionByName("main"));
}
TEST_F(ModuleTest, IsEntryPoint) {
type::F32Type f32;
Module m;
auto func = std::make_unique<Function>("other_func", VariableList{}, &f32);
m.AddFunction(std::move(func));
m.AddEntryPoint(
std::make_unique<EntryPoint>(PipelineStage::kVertex, "main", "vtx_main"));
EXPECT_TRUE(m.IsFunctionEntryPoint("vtx_main"));
EXPECT_FALSE(m.IsFunctionEntryPoint("other_func"));
}
TEST_F(ModuleTest, LookupFunctionMissing) {
Module m;
EXPECT_EQ(nullptr, m.FindFunctionByName("Missing"));
@ -155,25 +141,6 @@ TEST_F(ModuleTest, IsValid_Invalid_GlobalVariable) {
EXPECT_FALSE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_EntryPoint) {
Module m;
m.AddEntryPoint(
std::make_unique<EntryPoint>(PipelineStage::kVertex, "main", "vtx_main"));
EXPECT_TRUE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Null_EntryPoint) {
Module m;
m.AddEntryPoint(nullptr);
EXPECT_FALSE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Invalid_EntryPoint) {
Module m;
m.AddEntryPoint(std::make_unique<EntryPoint>());
EXPECT_FALSE(m.IsValid());
}
TEST_F(ModuleTest, IsValid_Alias) {
type::F32Type f32;
type::AliasType alias("alias", &f32);

22
src/ast/transform/vertex_pulling_transform.cc
View File

@ -77,29 +77,15 @@ bool VertexPullingTransform::Run() {
}
// Find entry point
EntryPoint* entry_point = nullptr;
for (const auto& entry : mod_->entry_points()) {
if (entry->name() == entry_point_name_ ||
(entry->name().empty() &&
entry->function_name() == entry_point_name_)) {
entry_point = entry.get();
break;
}
}
if (entry_point == nullptr) {
auto* func = mod_->FindFunctionByNameAndStage(entry_point_name_,
PipelineStage::kVertex);
if (func == nullptr) {
SetError("Vertex stage entry point not found");
return false;
}
// Check entry point is the right stage
if (entry_point->stage() != PipelineStage::kVertex) {
SetError("Entry point is not for vertex stage");
return false;
}
// Save the vertex function
auto* vertex_func = mod_->FindFunctionByName(entry_point->function_name());
auto* vertex_func = mod_->FindFunctionByName(func->name());
// TODO(idanr): Need to check shader locations in descriptor cover all
// attributes

56
src/ast/transform/vertex_pulling_transform_test.cc
View File

@ -17,6 +17,8 @@
#include "gtest/gtest.h"
#include "src/ast/decorated_variable.h"
#include "src/ast/function.h"
#include "src/ast/pipeline_stage.h"
#include "src/ast/stage_decoration.h"
#include "src/ast/type/array_type.h"
#include "src/ast/type/f32_type.h"
#include "src/ast/type/i32_type.h"
@ -38,11 +40,12 @@ class VertexPullingTransformHelper {
// Create basic module with an entry point and vertex function
void InitBasicModule() {
mod()->AddEntryPoint(std::make_unique<EntryPoint>(PipelineStage::kVertex,
"main", "vtx_main"));
mod()->AddFunction(std::make_unique<Function>(
"vtx_main", VariableList{},
ctx_.type_mgr().Get(std::make_unique<type::VoidType>())));
auto func = std::make_unique<Function>(
"main", VariableList{},
ctx_.type_mgr().Get(std::make_unique<type::VoidType>()));
func->add_decoration(
std::make_unique<ast::StageDecoration>(ast::PipelineStage ::kVertex));
mod()->AddFunction(std::move(func));
}
// Set up the transformation, after building the module
@ -71,6 +74,7 @@ class VertexPullingTransformHelper {
mod_->AddGlobalVariable(std::move(var));
}
Context* ctx() { return &ctx_; }
ast::Module* mod() { return mod_.get(); }
VertexPullingTransform* transform() { return transform_.get(); }
@ -104,12 +108,16 @@ TEST_F(VertexPullingTransformTest, Error_InvalidEntryPoint) {
}
TEST_F(VertexPullingTransformTest, Error_EntryPointWrongStage) {
InitBasicModule();
mod()->entry_points()[0]->set_pipeline_stage(PipelineStage::kFragment);
auto func = std::make_unique<Function>(
"main", VariableList{},
ctx()->type_mgr().Get(std::make_unique<type::VoidType>()));
func->add_decoration(
std::make_unique<ast::StageDecoration>(ast::PipelineStage::kFragment));
mod()->AddFunction(std::move(func));
InitTransform({});
EXPECT_FALSE(transform()->Run());
EXPECT_EQ(transform()->GetError(), "Entry point is not for vertex stage");
EXPECT_EQ(transform()->GetError(), "Vertex stage entry point not found");
}
TEST_F(VertexPullingTransformTest, BasicModule) {
@ -118,14 +126,6 @@ TEST_F(VertexPullingTransformTest, BasicModule) {
EXPECT_TRUE(transform()->Run());
}
TEST_F(VertexPullingTransformTest, EntryPointUsingFunctionName) {
InitBasicModule();
mod()->entry_points()[0]->set_name("");
InitTransform({});
transform()->SetEntryPoint("vtx_main");
EXPECT_TRUE(transform()->Run());
}
TEST_F(VertexPullingTransformTest, OneAttribute) {
InitBasicModule();
@ -159,8 +159,8 @@ TEST_F(VertexPullingTransformTest, OneAttribute) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{
@ -240,8 +240,8 @@ TEST_F(VertexPullingTransformTest, OneInstancedAttribute) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{
@ -321,8 +321,8 @@ TEST_F(VertexPullingTransformTest, OneAttributeDifferentOutputSet) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{
@ -454,8 +454,8 @@ TEST_F(VertexPullingTransformTest, ExistingVertexIndexAndInstanceIndex) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{
@ -573,8 +573,8 @@ TEST_F(VertexPullingTransformTest, TwoAttributesSameBuffer) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{
@ -777,8 +777,8 @@ TEST_F(VertexPullingTransformTest, FloatVectorAttributes) {
storage_buffer
__struct_
}
EntryPoint{vertex as main = vtx_main}
Function vtx_main -> __void
Function main -> __void
StageDecoration{vertex}
()
{
Block{

2
src/reader/wgsl/lexer.cc
View File

@ -503,8 +503,6 @@ Token Lexer::check_keyword(const Source& source, const std::string& str) {
return {Token::Type::kElse, source, "else"};
if (str == "elseif")
return {Token::Type::kElseIf, source, "elseif"};
if (str == "entry_point")
return {Token::Type::kEntryPoint, source, "entry_point"};
if (str == "f32")
return {Token::Type::kF32, source, "f32"};
if (str == "fallthrough")

1
src/reader/wgsl/lexer_test.cc
View File

@ -429,7 +429,6 @@ INSTANTIATE_TEST_SUITE_P(
TokenData{"discard", Token::Type::kDiscard},
TokenData{"else", Token::Type::kElse},
TokenData{"elseif", Token::Type::kElseIf},
TokenData{"entry_point", Token::Type::kEntryPoint},
TokenData{"f32", Token::Type::kF32},
TokenData{"fallthrough", Token::Type::kFallthrough},
TokenData{"false", Token::Type::kFalse},

68
src/reader/wgsl/parser_impl.cc
View File

@ -206,7 +206,6 @@ void ParserImpl::translation_unit() {
// | import_decl SEMICOLON
// | global_variable_decl SEMICLON
// | global_constant_decl SEMICOLON
// | entry_point_decl SEMICOLON
// | type_alias SEMICOLON
// | function_decl
void ParserImpl::global_decl() {
@ -258,19 +257,6 @@ void ParserImpl::global_decl() {
return;
}
auto ep = entry_point_decl();
if (has_error())
return;
if (ep != nullptr) {
t = next();
if (!t.IsSemicolon()) {
set_error(t, "missing ';' for entry point");
return;
}
module_.AddEntryPoint(std::move(ep));
return;
}
auto* ta = type_alias();
if (has_error())
return;
@ -2036,60 +2022,6 @@ ast::VariableList ParserImpl::param_list() {
return ret;
}
// entry_point_decl
// : ENTRY_POINT pipeline_stage EQUAL IDENT
// | ENTRY_POINT pipeline_stage AS STRING_LITERAL EQUAL IDENT
// | ENTRY_POINT pipeline_stage AS IDENT EQUAL IDENT
std::unique_ptr<ast::EntryPoint> ParserImpl::entry_point_decl() {
auto t = peek();
auto source = t.source();
if (!t.IsEntryPoint())
return nullptr;
next(); // Consume the peek
auto stage = pipeline_stage();
if (has_error())
return nullptr;
if (stage == ast::PipelineStage::kNone) {
set_error(peek(), "missing pipeline stage for entry point");
return nullptr;
}
t = next();
std::string name;
if (t.IsAs()) {
t = next();
if (t.IsStringLiteral()) {
name = t.to_str();
} else if (t.IsIdentifier()) {
name = t.to_str();
} else {
set_error(t, "invalid name for entry point");
return nullptr;
}
t = next();
}
if (!t.IsEqual()) {
set_error(t, "missing = for entry point");
return nullptr;
}
t = next();
if (!t.IsIdentifier()) {
set_error(t, "invalid function name for entry point");
return nullptr;
}
auto fn_name = t.to_str();
// Set the name to the function name if it isn't provided
if (name.length() == 0)
name = fn_name;
return std::make_unique<ast::EntryPoint>(source, stage, name, fn_name);
}
// pipeline_stage
// : VERTEX
// | FRAGMENT

4
src/reader/wgsl/parser_impl.h
View File

@ -27,7 +27,6 @@
#include "src/ast/case_statement.h"
#include "src/ast/constructor_expression.h"
#include "src/ast/else_statement.h"
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/import.h"
#include "src/ast/literal.h"
@ -222,9 +221,6 @@ class ParserImpl {
/// Parses a `param_list` grammar element
/// @returns the parsed variables
ast::VariableList param_list();
/// Parses a `entry_point_decl` grammar element
/// @returns the EntryPoint or nullptr on error
std::unique_ptr<ast::EntryPoint> entry_point_decl();
/// Parses a `pipeline_stage` grammar element
/// @returns the pipeline stage or PipelineStage::kNone if none matched
ast::PipelineStage pipeline_stage();

122
src/reader/wgsl/parser_impl_entry_point_decl_test.cc
View File

@ -1,122 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "gtest/gtest.h"
#include "src/ast/variable.h"
#include "src/reader/wgsl/parser_impl.h"
#include "src/reader/wgsl/parser_impl_test_helper.h"
namespace tint {
namespace reader {
namespace wgsl {
namespace {
TEST_F(ParserImplTest, EntryPoint_Parses) {
auto* p = parser("entry_point fragment = main");
auto e = p->entry_point_decl();
ASSERT_NE(e, nullptr);
ASSERT_FALSE(p->has_error());
EXPECT_EQ(e->stage(), ast::PipelineStage::kFragment);
EXPECT_EQ(e->name(), "main");
EXPECT_EQ(e->function_name(), "main");
}
TEST_F(ParserImplTest, EntryPoint_ParsesWithStringName) {
auto* p = parser(R"(entry_point vertex as "main" = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_NE(e, nullptr);
ASSERT_FALSE(p->has_error());
EXPECT_EQ(e->stage(), ast::PipelineStage::kVertex);
EXPECT_EQ(e->name(), "main");
EXPECT_EQ(e->function_name(), "vtx_main");
}
TEST_F(ParserImplTest, EntryPoint_ParsesWithIdentName) {
auto* p = parser(R"(entry_point vertex as main = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_NE(e, nullptr);
ASSERT_FALSE(p->has_error());
EXPECT_EQ(e->stage(), ast::PipelineStage::kVertex);
EXPECT_EQ(e->name(), "main");
EXPECT_EQ(e->function_name(), "vtx_main");
}
TEST_F(ParserImplTest, EntryPoint_MissingFnName) {
auto* p = parser(R"(entry_point vertex as main =)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:29: invalid function name for entry point");
}
TEST_F(ParserImplTest, EntryPoint_InvalidFnName) {
auto* p = parser(R"(entry_point vertex as main = 123)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:30: invalid function name for entry point");
}
TEST_F(ParserImplTest, EntryPoint_MissingEqual) {
auto* p = parser(R"(entry_point vertex as main vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:28: missing = for entry point");
}
TEST_F(ParserImplTest, EntryPoint_MissingName) {
auto* p = parser(R"(entry_point vertex as = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:23: invalid name for entry point");
}
TEST_F(ParserImplTest, EntryPoint_InvalidName) {
auto* p = parser(R"(entry_point vertex as 123 = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:23: invalid name for entry point");
}
TEST_F(ParserImplTest, EntryPoint_MissingStageWithIdent) {
auto* p = parser(R"(entry_point as 123 = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:13: missing pipeline stage for entry point");
}
TEST_F(ParserImplTest, EntryPoint_MissingStage) {
auto* p = parser(R"(entry_point = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:13: missing pipeline stage for entry point");
}
TEST_F(ParserImplTest, EntryPoint_InvalidStage) {
auto* p = parser(R"(entry_point invalid = vtx_main)");
auto e = p->entry_point_decl();
ASSERT_TRUE(p->has_error());
ASSERT_EQ(e, nullptr);
EXPECT_EQ(p->error(), "1:13: missing pipeline stage for entry point");
}
} // namespace
} // namespace wgsl
} // namespace reader
} // namespace tint

24
src/reader/wgsl/parser_impl_global_decl_test.cc
View File

@ -108,30 +108,6 @@ TEST_F(ParserImplTest, GlobalDecl_GlobalConstant_MissingSemicolon) {
EXPECT_EQ(p->error(), "1:38: missing ';' for constant declaration");
}
TEST_F(ParserImplTest, GlobalDecl_EntryPoint) {
auto* p = parser("entry_point vertex = main;");
p->global_decl();
ASSERT_FALSE(p->has_error()) << p->error();
auto m = p->module();
ASSERT_EQ(m.entry_points().size(), 1u);
EXPECT_EQ(m.entry_points()[0]->name(), "main");
}
TEST_F(ParserImplTest, GlobalDecl_EntryPoint_Invalid) {
auto* p = parser("entry_point main;");
p->global_decl();
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:13: missing pipeline stage for entry point");
}
TEST_F(ParserImplTest, GlobalDecl_EntryPoint_MissingSemicolon) {
auto* p = parser("entry_point vertex = main");
p->global_decl();
ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:26: missing ';' for entry point");
}
TEST_F(ParserImplTest, GlobalDecl_TypeAlias) {
auto* p = parser("type A = i32;");
p->global_decl();

3
src/reader/wgsl/parser_impl_test.cc
View File

@ -34,7 +34,7 @@ import "GLSL.std.430" as glsl;
[[location(0)]] var<out> gl_FragColor : vec4<f32>;
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
gl_FragColor = vec4<f32>(.4, .2, .3, 1);
}
@ -43,7 +43,6 @@ fn main() -> void {
auto m = p->module();
ASSERT_EQ(1u, m.imports().size());
ASSERT_EQ(1u, m.entry_points().size());
ASSERT_EQ(1u, m.functions().size());
ASSERT_EQ(1u, m.global_variables().size());
}

3
src/reader/wgsl/parser_test.cc
View File

@ -38,7 +38,7 @@ import "GLSL.std.430" as glsl;
[[location(0)]] var<out> gl_FragColor : vec4<f32>;
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
gl_FragColor = vec4<f32>(.4, .2, .3, 1);
}
@ -47,7 +47,6 @@ fn main() -> void {
auto m = p.module();
ASSERT_EQ(1u, m.imports().size());
ASSERT_EQ(1u, m.entry_points().size());
ASSERT_EQ(1u, m.functions().size());
ASSERT_EQ(1u, m.global_variables().size());
}

2
src/reader/wgsl/token.cc
View File

@ -139,8 +139,6 @@ std::string Token::TypeToName(Type type) {
return "else";
case Token::Type::kElseIf:
return "elseif";
case Token::Type::kEntryPoint:
return "entry_point";
case Token::Type::kF32:
return "f32";
case Token::Type::kFallthrough:

4
src/reader/wgsl/token.h
View File

@ -150,8 +150,6 @@ class Token {
kElse,
/// A 'elseif'
kElseIf,
/// A 'entry_point'
kEntryPoint,
/// A 'f32'
kF32,
/// A 'fallthrough'
@ -531,8 +529,6 @@ class Token {
bool IsElse() const { return type_ == Type::kElse; }
/// @returns true if token is a 'elseif'
bool IsElseIf() const { return type_ == Type::kElseIf; }
/// @returns true if token is a 'entry_point'
bool IsEntryPoint() const { return type_ == Type::kEntryPoint; }
/// @returns true if token is a 'f32'
bool IsF32() const { return type_ == Type::kF32; }
/// @returns true if token is a 'fallthrough'

8
src/type_determiner.cc
View File

@ -209,14 +209,6 @@ bool TypeDeterminer::Determine() {
return false;
}
// Walk over the caller to callee information and update functions with which
// entry points call those functions.
for (const auto& ep : mod_->entry_points()) {
for (const auto& callee : caller_to_callee_[ep->function_name()]) {
set_entry_points(callee, ep->name());
}
}
// Walk over the caller to callee information and update functions with which
// entry points call those functions.
for (const auto& func : mod_->functions()) {

108
src/type_determiner_test.cc
View File

@ -4373,114 +4373,6 @@ INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
testing::Values(GLSLData{"sclamp", GLSLstd450SClamp},
GLSLData{"uclamp", GLSLstd450UClamp}));
TEST_F(TypeDeterminerTest, Function_EntryPoints) {
ast::type::F32Type f32;
// fn b() {}
// fn c() { b(); }
// fn a() { c(); }
// fn ep_1() { a(); b(); }
// fn ep_2() { c();}
//
// c -> {ep_1, ep_2}
// a -> {ep_1}
// b -> {ep_1, ep_2}
// ep_1 -> {}
// ep_2 -> {}
ast::VariableList params;
auto func_b = std::make_unique<ast::Function>("b", std::move(params), &f32);
auto* func_b_ptr = func_b.get();
auto body = std::make_unique<ast::BlockStatement>();
func_b->set_body(std::move(body));
auto func_c = std::make_unique<ast::Function>("c", std::move(params), &f32);
auto* func_c_ptr = func_c.get();
body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("second"),
std::make_unique<ast::CallExpression>(
std::make_unique<ast::IdentifierExpression>("b"),
ast::ExpressionList{})));
func_c->set_body(std::move(body));
auto func_a = std::make_unique<ast::Function>("a", std::move(params), &f32);
auto* func_a_ptr = func_a.get();
body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("first"),
std::make_unique<ast::CallExpression>(
std::make_unique<ast::IdentifierExpression>("c"),
ast::ExpressionList{})));
func_a->set_body(std::move(body));
auto ep_1_func =
std::make_unique<ast::Function>("ep_1_func", std::move(params), &f32);
auto* ep_1_func_ptr = ep_1_func.get();
body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("call_a"),
std::make_unique<ast::CallExpression>(
std::make_unique<ast::IdentifierExpression>("a"),
ast::ExpressionList{})));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("call_b"),
std::make_unique<ast::CallExpression>(
std::make_unique<ast::IdentifierExpression>("b"),
ast::ExpressionList{})));
ep_1_func->set_body(std::move(body));
auto ep_2_func =
std::make_unique<ast::Function>("ep_2_func", std::move(params), &f32);
auto* ep_2_func_ptr = ep_2_func.get();
body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("call_c"),
std::make_unique<ast::CallExpression>(
std::make_unique<ast::IdentifierExpression>("c"),
ast::ExpressionList{})));
ep_2_func->set_body(std::move(body));
auto ep_1 = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex,
"ep_1", "ep_1_func");
auto ep_2 = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex,
"ep_2", "ep_2_func");
mod()->AddFunction(std::move(func_b));
mod()->AddFunction(std::move(func_c));
mod()->AddFunction(std::move(func_a));
mod()->AddFunction(std::move(ep_1_func));
mod()->AddFunction(std::move(ep_2_func));
mod()->AddEntryPoint(std::move(ep_1));
mod()->AddEntryPoint(std::move(ep_2));
// Register the functions and calculate the callers
ASSERT_TRUE(td()->Determine()) << td()->error();
const auto& b_eps = func_b_ptr->ancestor_entry_points();
ASSERT_EQ(2u, b_eps.size());
EXPECT_EQ("ep_1", b_eps[0]);
EXPECT_EQ("ep_2", b_eps[1]);
const auto& a_eps = func_a_ptr->ancestor_entry_points();
ASSERT_EQ(1u, a_eps.size());
EXPECT_EQ("ep_1", a_eps[0]);
const auto& c_eps = func_c_ptr->ancestor_entry_points();
ASSERT_EQ(2u, c_eps.size());
EXPECT_EQ("ep_1", c_eps[0]);
EXPECT_EQ("ep_2", c_eps[1]);
EXPECT_TRUE(ep_1_func_ptr->ancestor_entry_points().empty());
EXPECT_TRUE(ep_2_func_ptr->ancestor_entry_points().empty());
}
TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) {
ast::type::F32Type f32;

4
src/validator_function_test.cc
View File

@ -17,7 +17,6 @@
#include "gtest/gtest.h"
#include "spirv/unified1/GLSL.std.450.h"
#include "src/ast/call_statement.h"
#include "src/ast/entry_point.h"
#include "src/ast/pipeline_stage.h"
#include "src/ast/return_statement.h"
#include "src/ast/scalar_constructor_expression.h"
@ -91,8 +90,7 @@ TEST_F(ValidateFunctionTest, FunctionTypeMustMatchReturnStatementType_Pass) {
mod()->AddFunction(std::move(func));
EXPECT_TRUE(td()->DetermineFunctions(mod()->functions())) << td()->error();
EXPECT_TRUE(v()->ValidateFunctions(mod(), mod()->functions()))
<< v()->error();
EXPECT_TRUE(v()->ValidateFunctions(mod()->functions())) << v()->error();
}
TEST_F(ValidateFunctionTest, FunctionTypeMustMatchReturnStatementType_fail) {

7
src/validator_impl.cc
View File

@ -49,7 +49,7 @@ bool ValidatorImpl::Validate(const ast::Module* module) {
if (!CheckImports(module)) {
return false;
}
if (!ValidateFunctions(module, module->functions())) {
if (!ValidateFunctions(module->functions())) {
return false;
}
@ -82,8 +82,7 @@ bool ValidatorImpl::ValidateGlobalVariables(
return true;
}
bool ValidatorImpl::ValidateFunctions(const ast::Module* mod,
const ast::FunctionList& funcs) {
bool ValidatorImpl::ValidateFunctions(const ast::FunctionList& funcs) {
ScopeStack<ast::PipelineStage> entry_point_map;
entry_point_map.push_scope();
@ -135,7 +134,7 @@ bool ValidatorImpl::ValidateFunctions(const ast::Module* mod,
current_function_ = nullptr;
}
if (pipeline_count == 0 && mod->entry_points().empty()) {
if (pipeline_count == 0) {
set_error(Source{0, 0},
"v-0003: At least one of vertex, fragment or compute shader must "
"be present");

10
src/validator_impl.h
View File

@ -20,7 +20,6 @@
#include "src/ast/assignment_statement.h"
#include "src/ast/call_expression.h"
#include "src/ast/entry_point.h"
#include "src/ast/expression.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/module.h"
@ -58,11 +57,9 @@ class ValidatorImpl {
/// @returns true if the validation was successful
bool ValidateGlobalVariables(const ast::VariableList& global_vars);
/// Validates Functions
/// @param mod the module
/// @param funcs the functions to check
/// @returns true if the validation was successful
bool ValidateFunctions(const ast::Module* mod,
const ast::FunctionList& funcs);
bool ValidateFunctions(const ast::FunctionList& funcs);
/// Validates a function
/// @param func the function to check
/// @returns true if the validation was successful
@ -112,11 +109,6 @@ class ValidatorImpl {
/// @param expr the call to validate
/// @returns true if successful
bool ValidateCallExpr(const ast::CallExpression* expr);
/// Validates entry points
/// this funtion must be called after populating function_stack_
/// @param eps the vector of entry points to check
/// @return true if the validation was successful
bool ValidateEntryPoints(const ast::EntryPointList& eps);
/// Validates switch statements
/// @param s the switch statement to check
/// @returns true if the valdiation was successful

273
src/writer/hlsl/generator_impl.cc
View File

@ -134,12 +134,6 @@ bool GeneratorImpl::Generate(std::ostream& out) {
}
}
for (const auto& ep : module_->entry_points()) {
if (!EmitEntryPointData(out, ep.get())) {
return false;
}
}
// Make sure all entry point data is emitted before the entry point functions
for (const auto& func : module_->functions()) {
if (!func->IsEntryPoint()) {
@ -156,12 +150,6 @@ bool GeneratorImpl::Generate(std::ostream& out) {
return false;
}
}
for (const auto& ep : module_->entry_points()) {
if (!EmitEntryPointFunction(out, ep.get())) {
return false;
}
out << std::endl;
}
for (const auto& func : module_->functions()) {
if (!func->IsEntryPoint()) {
@ -1106,7 +1094,7 @@ bool GeneratorImpl::EmitFunction(std::ostream& out, ast::Function* func) {
make_indent(out);
// Entry points will be emitted later, skip for now.
if (func->IsEntryPoint() || module_->IsFunctionEntryPoint(func->name())) {
if (func->IsEntryPoint()) {
return true;
}
@ -1208,200 +1196,6 @@ bool GeneratorImpl::EmitFunctionInternal(std::ostream& out,
return true;
}
bool GeneratorImpl::EmitEntryPointData(std::ostream& out, ast::EntryPoint* ep) {
auto* func = module_->FindFunctionByName(ep->function_name());
if (func == nullptr) {
error_ = "Unable to find entry point function: " + ep->function_name();
return false;
}
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()) {
auto* var = data.first;
auto* deco = data.second;
if (var->storage_class() == ast::StorageClass::kInput) {
in_variables.push_back({var, deco});
} else if (var->storage_class() == ast::StorageClass::kOutput) {
outvariables.push_back({var, deco});
}
}
for (auto data : func->referenced_builtin_variables()) {
auto* var = data.first;
auto* deco = data.second;
if (var->storage_class() == ast::StorageClass::kInput) {
in_variables.push_back({var, deco});
} else if (var->storage_class() == ast::StorageClass::kOutput) {
outvariables.push_back({var, deco});
}
}
bool emitted_uniform = false;
for (auto data : func->referenced_uniform_variables()) {
auto* var = data.first;
// TODO(dsinclair): We're using the binding to make up the buffer number but
// we should instead be using a provided mapping that uses both buffer and
// set. https://bugs.chromium.org/p/tint/issues/detail?id=104
auto* binding = data.second.binding;
if (binding == nullptr) {
error_ = "unable to find binding information for uniform: " + var->name();
return false;
}
// auto* set = data.second.set;
auto* type = var->type()->UnwrapAliasesIfNeeded();
if (type->IsStruct()) {
auto* strct = type->AsStruct();
out << "ConstantBuffer<" << strct->name() << "> " << var->name()
<< " : register(b" << binding->value() << ");" << std::endl;
} else {
// TODO(dsinclair): There is outstanding spec work to require all uniform
// buffers to be [[block]] decorated, which means structs. This is
// currently not the case, so this code handles the cases where the data
// is not a block.
// Relevant: https://github.com/gpuweb/gpuweb/issues/1004
// https://github.com/gpuweb/gpuweb/issues/1008
out << "cbuffer : register(b" << binding->value() << ") {" << std::endl;
increment_indent();
make_indent(out);
if (!EmitType(out, type, "")) {
return false;
}
out << " " << var->name() << ";" << std::endl;
decrement_indent();
out << "};" << std::endl;
}
emitted_uniform = true;
}
if (emitted_uniform) {
out << std::endl;
}
bool emitted_storagebuffer = false;
for (auto data : func->referenced_storagebuffer_variables()) {
auto* var = data.first;
auto* binding = data.second.binding;
out << "RWByteAddressBuffer " << var->name() << " : register(u"
<< binding->value() << ");" << std::endl;
emitted_storagebuffer = true;
}
if (emitted_storagebuffer) {
out << std::endl;
}
auto ep_name = ep->name();
if (ep_name.empty()) {
ep_name = ep->function_name();
}
// TODO(dsinclair): There is a potential bug here. Entry points can have the
// same name in WGSL if they have different pipeline stages. This does not
// take that into account and will emit duplicate struct names. I'm ignoring
// this until https://github.com/gpuweb/gpuweb/issues/662 is resolved as it
// may remove this issue and entry point names will need to be unique.
if (!in_variables.empty()) {
auto in_struct_name = generate_name(ep_name + "_" + kInStructNameSuffix);
auto in_var_name = generate_name(kTintStructInVarPrefix);
ep_name_to_in_data_[ep_name] = {in_struct_name, in_var_name};
make_indent(out);
out << "struct " << in_struct_name << " {" << std::endl;
increment_indent();
for (auto& data : in_variables) {
auto* var = data.first;
auto* deco = data.second;
make_indent(out);
if (!EmitType(out, var->type(), var->name())) {
return false;
}
out << " " << var->name() << " : ";
if (deco->IsLocation()) {
if (ep->stage() == ast::PipelineStage::kCompute) {
error_ = "invalid location variable for pipeline stage";
return false;
}
out << "TEXCOORD" << deco->AsLocation()->value();
} else if (deco->IsBuiltin()) {
auto attr = builtin_to_attribute(deco->AsBuiltin()->value());
if (attr.empty()) {
error_ = "unsupported builtin";
return false;
}
out << attr;
} else {
error_ = "unsupported variable decoration for entry point output";
return false;
}
out << ";" << std::endl;
}
decrement_indent();
make_indent(out);
out << "};" << std::endl << std::endl;
}
if (!outvariables.empty()) {
auto outstruct_name = generate_name(ep_name + "_" + kOutStructNameSuffix);
auto outvar_name = generate_name(kTintStructOutVarPrefix);
ep_name_to_out_data_[ep_name] = {outstruct_name, outvar_name};
make_indent(out);
out << "struct " << outstruct_name << " {" << std::endl;
increment_indent();
for (auto& data : outvariables) {
auto* var = data.first;
auto* deco = data.second;
make_indent(out);
if (!EmitType(out, var->type(), var->name())) {
return false;
}
out << " " << var->name() << " : ";
if (deco->IsLocation()) {
auto loc = deco->AsLocation()->value();
if (ep->stage() == ast::PipelineStage::kVertex) {
out << "TEXCOORD" << loc;
} else if (ep->stage() == ast::PipelineStage::kFragment) {
out << "SV_Target" << loc << "";
} else {
error_ = "invalid location variable for pipeline stage";
return false;
}
} else if (deco->IsBuiltin()) {
auto attr = builtin_to_attribute(deco->AsBuiltin()->value());
if (attr.empty()) {
error_ = "unsupported builtin";
return false;
}
out << attr;
} else {
error_ = "unsupported variable decoration for entry point output";
return false;
}
out << ";" << std::endl;
}
decrement_indent();
make_indent(out);
out << "};" << std::endl << std::endl;
}
return true;
}
bool GeneratorImpl::EmitEntryPointData(std::ostream& out, ast::Function* func) {
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> in_variables;
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>> outvariables;
@ -1608,71 +1402,6 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const {
return "";
}
bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out,
ast::EntryPoint* ep) {
make_indent(out);
current_ep_name_ = ep->name();
if (current_ep_name_.empty()) {
current_ep_name_ = ep->function_name();
}
auto* func = module_->FindFunctionByName(ep->function_name());
if (func == nullptr) {
error_ = "unable to find function for entry point: " + ep->function_name();
return false;
}
if (ep->stage() == ast::PipelineStage::kCompute) {
uint32_t x = 0;
uint32_t y = 0;
uint32_t z = 0;
std::tie(x, y, z) = func->workgroup_size();
out << "[numthreads(" << std::to_string(x) << ", " << std::to_string(y)
<< ", " << std::to_string(z) << ")]" << std::endl;
make_indent(out);
}
auto outdata = ep_name_to_out_data_.find(current_ep_name_);
bool has_outdata = outdata != ep_name_to_out_data_.end();
if (has_outdata) {
out << outdata->second.struct_name;
} else {
out << "void";
}
out << " " << namer_.NameFor(current_ep_name_) << "(";
auto in_data = ep_name_to_in_data_.find(current_ep_name_);
if (in_data != ep_name_to_in_data_.end()) {
out << in_data->second.struct_name << " " << in_data->second.var_name;
}
out << ") {" << std::endl;
increment_indent();
if (has_outdata) {
make_indent(out);
out << outdata->second.struct_name << " " << outdata->second.var_name << ";"
<< std::endl;
}
generating_entry_point_ = true;
for (const auto& s : *(func->body())) {
if (!EmitStatement(out, s.get())) {
return false;
}
}
generating_entry_point_ = false;
decrement_indent();
make_indent(out);
out << "}" << std::endl;
current_ep_name_ = "";
return true;
}
bool GeneratorImpl::EmitEntryPointFunction(std::ostream& out,
ast::Function* func) {
make_indent(out);

10
src/writer/hlsl/generator_impl.h
View File

@ -193,21 +193,11 @@ class GeneratorImpl {
const std::string& ep_name);
/// Handles emitting information for an entry point
/// @param out the output stream
/// @param ep the entry point
/// @returns true if the entry point data was emitted
bool EmitEntryPointData(std::ostream& out, ast::EntryPoint* ep);
/// Handles emitting information for an entry point
/// @param out the output stream
/// @param func the entry point
/// @returns true if the entry point data was emitted
bool EmitEntryPointData(std::ostream& out, ast::Function* func);
/// Handles emitting the entry point function
/// @param out the output stream
/// @param ep the entry point
/// @returns true if the entry point function was emitted
bool EmitEntryPointFunction(std::ostream& out, ast::EntryPoint* ep);
/// Handles emitting the entry point function
/// @param out the output stream
/// @param func the entry point
/// @returns true if the entry point function was emitted
bool EmitEntryPointFunction(std::ostream& out, ast::Function* func);

463
src/writer/hlsl/generator_impl_entry_point_test.cc
View File

@ -1,463 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/ast/assignment_statement.h"
#include "src/ast/decorated_variable.h"
#include "src/ast/entry_point.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/location_decoration.h"
#include "src/ast/member_accessor_expression.h"
#include "src/ast/module.h"
#include "src/ast/return_statement.h"
#include "src/ast/type/f32_type.h"
#include "src/ast/type/i32_type.h"
#include "src/ast/type/vector_type.h"
#include "src/ast/type/void_type.h"
#include "src/ast/variable.h"
#include "src/context.h"
#include "src/type_determiner.h"
#include "src/writer/hlsl/test_helper.h"
namespace tint {
namespace writer {
namespace hlsl {
namespace {
using HlslGeneratorImplTest_EntryPoint = TestHelper;
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Vertex_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// struct vtx_main_in {
// float foo : TEXCOORD0;
// int bar : TEXCOORD1;
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("vtx_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex, "",
"vtx_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(result(), R"(struct vtx_main_in {
float foo : TEXCOORD0;
int bar : TEXCOORD1;
};
)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Vertex_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// struct vtx_main_out {
// float foo : TEXCOORD0;
// int bar : TEXCOORD1;
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("vtx_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex, "",
"vtx_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(result(), R"(struct vtx_main_out {
float foo : TEXCOORD0;
int bar : TEXCOORD1;
};
)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Fragment_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// struct frag_main_in {
// float foo : TEXCOORD0;
// int bar : TEXCOORD1;
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("frag_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(result(), R"(struct main_in {
float foo : TEXCOORD0;
int bar : TEXCOORD1;
};
)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Fragment_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// struct frag_main_out {
// float foo : SV_Target0;
// int bar : SV_Target1;
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("frag_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(result(), R"(struct main_out {
float foo : SV_Target0;
int bar : SV_Target1;
};
)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Compute_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// -> Error, not allowed
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("comp_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kCompute,
"main", "comp_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_FALSE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(gen().error(), R"(invalid location variable for pipeline stage)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Compute_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// -> Error not allowed
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(foo_var.get());
td().RegisterVariableForTesting(bar_var.get());
mod()->AddGlobalVariable(std::move(foo_var));
mod()->AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("comp_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kCompute,
"main", "comp_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_FALSE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(gen().error(), R"(invalid location variable for pipeline stage)");
}
TEST_F(HlslGeneratorImplTest_EntryPoint, EmitEntryPointData_Builtins) {
// [[builtin(frag_coord)]] var<in> coord : vec4<f32>;
// [[builtin(frag_depth)]] var<out> depth : f32;
//
// struct main_in {
// vector<float, 4> coord : SV_Position;
// };
//
// struct main_out {
// float depth : SV_Depth;
// };
ast::type::F32Type f32;
ast::type::VoidType void_type;
ast::type::VectorType vec4(&f32, 4);
auto coord_var =
std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
"coord", ast::StorageClass::kInput, &vec4));
ast::VariableDecorationList decos;
decos.push_back(
std::make_unique<ast::BuiltinDecoration>(ast::Builtin::kFragCoord));
coord_var->set_decorations(std::move(decos));
auto depth_var =
std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
"depth", ast::StorageClass::kOutput, &f32));
decos.push_back(
std::make_unique<ast::BuiltinDecoration>(ast::Builtin::kFragDepth));
depth_var->set_decorations(std::move(decos));
td().RegisterVariableForTesting(coord_var.get());
td().RegisterVariableForTesting(depth_var.get());
mod()->AddGlobalVariable(std::move(coord_var));
mod()->AddGlobalVariable(std::move(depth_var));
ast::VariableList params;
auto func = std::make_unique<ast::Function>("frag_main", std::move(params),
&void_type);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("depth"),
std::make_unique<ast::MemberAccessorExpression>(
std::make_unique<ast::IdentifierExpression>("coord"),
std::make_unique<ast::IdentifierExpression>("x"))));
func->set_body(std::move(body));
mod()->AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod()->AddEntryPoint(std::move(ep));
ASSERT_TRUE(td().Determine()) << td().error();
ASSERT_TRUE(gen().EmitEntryPointData(out(), ep_ptr)) << gen().error();
EXPECT_EQ(result(), R"(struct main_in {
vector<float, 4> coord : SV_Position;
};
struct main_out {
float depth : SV_Depth;
};
)");
}
} // namespace
} // namespace hlsl
} // namespace writer
} // namespace tint

1
src/writer/hlsl/generator_impl_function_entry_point_data_test.cc
View File

@ -14,7 +14,6 @@
#include "src/ast/assignment_statement.h"
#include "src/ast/decorated_variable.h"
#include "src/ast/entry_point.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/location_decoration.h"
#include "src/ast/member_accessor_expression.h"

1132
src/writer/hlsl/generator_impl_function_test.cc
View File

File diff suppressed because it is too large Load Diff

8
src/writer/hlsl/generator_impl_test.cc
View File

@ -14,7 +14,6 @@
#include <memory>
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/module.h"
@ -30,10 +29,9 @@ using HlslGeneratorImplTest = TestHelper;
TEST_F(HlslGeneratorImplTest, Generate) {
ast::type::VoidType void_type;
mod()->AddFunction(std::make_unique<ast::Function>(
"my_func", ast::VariableList{}, &void_type));
mod()->AddEntryPoint(std::make_unique<ast::EntryPoint>(
ast::PipelineStage::kFragment, "", "my_func"));
auto func = std::make_unique<ast::Function>("my_func", ast::VariableList{},
&void_type);
mod()->AddFunction(std::move(func));
ASSERT_TRUE(gen().Generate(out())) << gen().error();
EXPECT_EQ(result(), R"(void my_func() {

276
src/writer/msl/generator_impl.cc
View File

@ -124,12 +124,6 @@ bool GeneratorImpl::Generate() {
}
}
for (const auto& ep : module_->entry_points()) {
if (!EmitEntryPointData(ep.get())) {
return false;
}
}
// Make sure all entry point data is emitted before the entry point functions
for (const auto& func : module_->functions()) {
if (!func->IsEntryPoint()) {
@ -147,12 +141,6 @@ bool GeneratorImpl::Generate() {
}
}
for (const auto& ep : module_->entry_points()) {
if (!EmitEntryPointFunction(ep.get())) {
return false;
}
out_ << std::endl;
}
for (const auto& func : module_->functions()) {
if (!func->IsEntryPoint()) {
continue;
@ -904,133 +892,6 @@ bool GeneratorImpl::EmitLiteral(ast::Literal* lit) {
return true;
}
bool GeneratorImpl::EmitEntryPointData(ast::EntryPoint* ep) {
auto* func = module_->FindFunctionByName(ep->function_name());
if (func == nullptr) {
error_ = "Unable to find entry point function: " + ep->function_name();
return false;
}
std::vector<std::pair<ast::Variable*, uint32_t>> in_locations;
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>>
out_variables;
for (auto data : func->referenced_location_variables()) {
auto* var = data.first;
auto* deco = data.second;
if (var->storage_class() == ast::StorageClass::kInput) {
in_locations.push_back({var, deco->value()});
} else if (var->storage_class() == ast::StorageClass::kOutput) {
out_variables.push_back({var, deco});
}
}
for (auto data : func->referenced_builtin_variables()) {
auto* var = data.first;
auto* deco = data.second;
if (var->storage_class() == ast::StorageClass::kOutput) {
out_variables.push_back({var, deco});
}
}
auto ep_name = ep->name();
if (ep_name.empty()) {
ep_name = ep->function_name();
}
// TODO(dsinclair): There is a potential bug here. Entry points can have the
// same name in WGSL if they have different pipeline stages. This does not
// take that into account and will emit duplicate struct names. I'm ignoring
// this until https://github.com/gpuweb/gpuweb/issues/662 is resolved as it
// may remove this issue and entry point names will need to be unique.
if (!in_locations.empty()) {
auto in_struct_name = generate_name(ep_name + "_" + kInStructNameSuffix);
auto in_var_name = generate_name(kTintStructInVarPrefix);
ep_name_to_in_data_[ep_name] = {in_struct_name, in_var_name};
make_indent();
out_ << "struct " << in_struct_name << " {" << std::endl;
increment_indent();
for (auto& data : in_locations) {
auto* var = data.first;
uint32_t loc = data.second;
make_indent();
if (!EmitType(var->type(), var->name())) {
return false;
}
out_ << " " << var->name() << " [[";
if (ep->stage() == ast::PipelineStage::kVertex) {
out_ << "attribute(" << loc << ")";
} else if (ep->stage() == ast::PipelineStage::kFragment) {
out_ << "user(locn" << loc << ")";
} else {
error_ = "invalid location variable for pipeline stage";
return false;
}
out_ << "]];" << std::endl;
}
decrement_indent();
make_indent();
out_ << "};" << std::endl << std::endl;
}
if (!out_variables.empty()) {
auto out_struct_name = generate_name(ep_name + "_" + kOutStructNameSuffix);
auto out_var_name = generate_name(kTintStructOutVarPrefix);
ep_name_to_out_data_[ep_name] = {out_struct_name, out_var_name};
make_indent();
out_ << "struct " << out_struct_name << " {" << std::endl;
increment_indent();
for (auto& data : out_variables) {
auto* var = data.first;
auto* deco = data.second;
make_indent();
if (!EmitType(var->type(), var->name())) {
return false;
}
out_ << " " << var->name() << " [[";
if (deco->IsLocation()) {
auto loc = deco->AsLocation()->value();
if (ep->stage() == ast::PipelineStage::kVertex) {
out_ << "user(locn" << loc << ")";
} else if (ep->stage() == ast::PipelineStage::kFragment) {
out_ << "color(" << loc << ")";
} else {
error_ = "invalid location variable for pipeline stage";
return false;
}
} else if (deco->IsBuiltin()) {
auto attr = builtin_to_attribute(deco->AsBuiltin()->value());
if (attr.empty()) {
error_ = "unsupported builtin";
return false;
}
out_ << attr;
} else {
error_ = "unsupported variable decoration for entry point output";
return false;
}
out_ << "]];" << std::endl;
}
decrement_indent();
make_indent();
out_ << "};" << std::endl << std::endl;
}
return true;
}
bool GeneratorImpl::EmitEntryPointData(ast::Function* func) {
std::vector<std::pair<ast::Variable*, uint32_t>> in_locations;
std::vector<std::pair<ast::Variable*, ast::VariableDecoration*>>
@ -1230,7 +1091,7 @@ bool GeneratorImpl::EmitFunction(ast::Function* func) {
make_indent();
// Entry points will be emitted later, skip for now.
if (func->IsEntryPoint() || module_->IsFunctionEntryPoint(func->name())) {
if (func->IsEntryPoint()) {
return true;
}
@ -1402,141 +1263,6 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const {
return "";
}
bool GeneratorImpl::EmitEntryPointFunction(ast::EntryPoint* ep) {
make_indent();
current_ep_name_ = ep->name();
if (current_ep_name_.empty()) {
current_ep_name_ = ep->function_name();
}
auto* func = module_->FindFunctionByName(ep->function_name());
if (func == nullptr) {
error_ = "unable to find function for entry point: " + ep->function_name();
return false;
}
EmitStage(ep->stage());
out_ << " ";
// This is an entry point, the return type is the entry point output structure
// if one exists, or void otherwise.
auto out_data = ep_name_to_out_data_.find(current_ep_name_);
bool has_out_data = out_data != ep_name_to_out_data_.end();
if (has_out_data) {
out_ << out_data->second.struct_name;
} else {
out_ << "void";
}
out_ << " " << namer_.NameFor(current_ep_name_) << "(";
bool first = true;
auto in_data = ep_name_to_in_data_.find(current_ep_name_);
if (in_data != ep_name_to_in_data_.end()) {
out_ << in_data->second.struct_name << " " << in_data->second.var_name
<< " [[stage_in]]";
first = false;
}
for (auto data : func->referenced_builtin_variables()) {
auto* var = data.first;
if (var->storage_class() != ast::StorageClass::kInput) {
continue;
}
if (!first) {
out_ << ", ";
}
first = false;
auto* builtin = data.second;
if (!EmitType(var->type(), "")) {
return false;
}
auto attr = builtin_to_attribute(builtin->value());
if (attr.empty()) {
error_ = "unknown builtin";
return false;
}
out_ << " " << var->name() << " [[" << attr << "]]";
}
for (auto data : func->referenced_uniform_variables()) {
if (!first) {
out_ << ", ";
}
first = false;
auto* var = data.first;
// TODO(dsinclair): We're using the binding to make up the buffer number but
// we should instead be using a provided mapping that uses both buffer and
// set. https://bugs.chromium.org/p/tint/issues/detail?id=104
auto* binding = data.second.binding;
if (binding == nullptr) {
error_ = "unable to find binding information for uniform: " + var->name();
return false;
}
// auto* set = data.second.set;
out_ << "constant ";
// TODO(dsinclair): Can you have a uniform array? If so, this needs to be
// updated to handle arrays property.
if (!EmitType(var->type(), "")) {
return false;
}
out_ << "& " << var->name() << " [[buffer(" << binding->value() << ")]]";
}
for (auto data : func->referenced_storagebuffer_variables()) {
if (!first) {
out_ << ", ";
}
first = false;
auto* var = data.first;
// TODO(dsinclair): We're using the binding to make up the buffer number but
// we should instead be using a provided mapping that uses both buffer and
// set. https://bugs.chromium.org/p/tint/issues/detail?id=104
auto* binding = data.second.binding;
// auto* set = data.second.set;
out_ << "device ";
// TODO(dsinclair): Can you have a storagebuffer have an array? If so, this
// needs to be updated to handle arrays property.
if (!EmitType(var->type(), "")) {
return false;
}
out_ << "& " << var->name() << " [[buffer(" << binding->value() << ")]]";
}
out_ << ") {" << std::endl;
increment_indent();
if (has_out_data) {
make_indent();
out_ << out_data->second.struct_name << " " << out_data->second.var_name
<< " = {};" << std::endl;
}
generating_entry_point_ = true;
for (const auto& s : *(func->body())) {
if (!EmitStatement(s.get())) {
return false;
}
}
generating_entry_point_ = false;
decrement_indent();
make_indent();
out_ << "}" << std::endl;
current_ep_name_ = "";
return true;
}
bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
make_indent();

8
src/writer/msl/generator_impl.h
View File

@ -118,18 +118,10 @@ class GeneratorImpl : public TextGenerator {
/// @returns true if the statement was emitted
bool EmitElse(ast::ElseStatement* stmt);
/// Handles emitting information for an entry point
/// @param ep the entry point
/// @returns true if the entry point data was emitted
bool EmitEntryPointData(ast::EntryPoint* ep);
/// Handles emitting information for an entry point
/// @param func the entry point function
/// @returns true if the entry point data was emitted
bool EmitEntryPointData(ast::Function* func);
/// Handles emitting the entry point function
/// @param ep the entry point
/// @returns true if the entry point function was emitted
bool EmitEntryPointFunction(ast::EntryPoint* ep);
/// Handles emitting the entry point function
/// @param func the entry point function
/// @returns true if the entry point function was emitted
bool EmitEntryPointFunction(ast::Function* func);

493
src/writer/msl/generator_impl_entry_point_test.cc
View File

@ -1,493 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "gtest/gtest.h"
#include "src/ast/assignment_statement.h"
#include "src/ast/decorated_variable.h"
#include "src/ast/entry_point.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/location_decoration.h"
#include "src/ast/member_accessor_expression.h"
#include "src/ast/module.h"
#include "src/ast/type/f32_type.h"
#include "src/ast/type/i32_type.h"
#include "src/ast/type/vector_type.h"
#include "src/ast/type/void_type.h"
#include "src/ast/variable.h"
#include "src/context.h"
#include "src/type_determiner.h"
#include "src/writer/msl/generator_impl.h"
namespace tint {
namespace writer {
namespace msl {
namespace {
using MslGeneratorImplTest = testing::Test;
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Vertex_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// struct vtx_main_in {
// float foo [[attribute(0)]];
// int bar [[attribute(1)]];
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("vtx_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex, "",
"vtx_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.result(), R"(struct vtx_main_in {
float foo [[attribute(0)]];
int bar [[attribute(1)]];
};
)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Vertex_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// struct vtx_main_out {
// float foo [[user(locn0)]];
// int bar [[user(locn1)]];
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("vtx_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kVertex, "",
"vtx_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.result(), R"(struct vtx_main_out {
float foo [[user(locn0)]];
int bar [[user(locn1)]];
};
)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Fragment_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// struct frag_main_in {
// float foo [[user(locn0)]];
// int bar [[user(locn1)]];
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("frag_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.result(), R"(struct main_in {
float foo [[user(locn0)]];
int bar [[user(locn1)]];
};
)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Fragment_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// struct frag_main_out {
// float foo [[color(0)]];
// int bar [[color(1)]];
// };
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("frag_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.result(), R"(struct main_out {
float foo [[color(0)]];
int bar [[color(1)]];
};
)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Compute_Input) {
// [[location(0)]] var<in> foo : f32;
// [[location(1)]] var<in> bar : i32;
//
// -> Error, not allowed
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kInput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kInput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("comp_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kCompute,
"main", "comp_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_FALSE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.error(), R"(invalid location variable for pipeline stage)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Compute_Output) {
// [[location(0)]] var<out> foo : f32;
// [[location(1)]] var<out> bar : i32;
//
// -> Error not allowed
ast::type::F32Type f32;
ast::type::I32Type i32;
auto foo_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("foo", ast::StorageClass::kOutput, &f32));
ast::VariableDecorationList decos;
decos.push_back(std::make_unique<ast::LocationDecoration>(0));
foo_var->set_decorations(std::move(decos));
auto bar_var = std::make_unique<ast::DecoratedVariable>(
std::make_unique<ast::Variable>("bar", ast::StorageClass::kOutput, &i32));
decos.push_back(std::make_unique<ast::LocationDecoration>(1));
bar_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(foo_var.get());
td.RegisterVariableForTesting(bar_var.get());
mod.AddGlobalVariable(std::move(foo_var));
mod.AddGlobalVariable(std::move(bar_var));
ast::VariableList params;
auto func =
std::make_unique<ast::Function>("comp_main", std::move(params), &f32);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("foo"),
std::make_unique<ast::IdentifierExpression>("foo")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("bar"),
std::make_unique<ast::IdentifierExpression>("bar")));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kCompute,
"main", "comp_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_FALSE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.error(), R"(invalid location variable for pipeline stage)");
}
TEST_F(MslGeneratorImplTest, EmitEntryPointData_Builtins) {
// Output builtins go in the output struct, input builtins will be passed
// as input parameters to the entry point function.
// [[builtin(frag_coord)]] var<in> coord : vec4<f32>;
// [[builtin(frag_depth)]] var<out> depth : f32;
//
// struct main_out {
// float depth [[depth(any)]];
// };
ast::type::F32Type f32;
ast::type::VoidType void_type;
ast::type::VectorType vec4(&f32, 4);
auto coord_var =
std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
"coord", ast::StorageClass::kInput, &vec4));
ast::VariableDecorationList decos;
decos.push_back(
std::make_unique<ast::BuiltinDecoration>(ast::Builtin::kFragCoord));
coord_var->set_decorations(std::move(decos));
auto depth_var =
std::make_unique<ast::DecoratedVariable>(std::make_unique<ast::Variable>(
"depth", ast::StorageClass::kOutput, &f32));
decos.push_back(
std::make_unique<ast::BuiltinDecoration>(ast::Builtin::kFragDepth));
depth_var->set_decorations(std::move(decos));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(coord_var.get());
td.RegisterVariableForTesting(depth_var.get());
mod.AddGlobalVariable(std::move(coord_var));
mod.AddGlobalVariable(std::move(depth_var));
ast::VariableList params;
auto func = std::make_unique<ast::Function>("frag_main", std::move(params),
&void_type);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("depth"),
std::make_unique<ast::MemberAccessorExpression>(
std::make_unique<ast::IdentifierExpression>("coord"),
std::make_unique<ast::IdentifierExpression>("x"))));
func->set_body(std::move(body));
mod.AddFunction(std::move(func));
auto ep = std::make_unique<ast::EntryPoint>(ast::PipelineStage::kFragment,
"main", "frag_main");
auto* ep_ptr = ep.get();
mod.AddEntryPoint(std::move(ep));
ASSERT_TRUE(td.Determine()) << td.error();
GeneratorImpl g(&mod);
ASSERT_TRUE(g.EmitEntryPointData(ep_ptr)) << g.error();
EXPECT_EQ(g.result(), R"(struct main_out {
float depth [[depth(any)]];
};
)");
}
} // namespace
} // namespace msl
} // namespace writer
} // namespace tint

992
src/writer/msl/generator_impl_function_test.cc
View File

File diff suppressed because it is too large Load Diff

12
src/writer/msl/generator_impl_test.cc
View File

@ -17,11 +17,11 @@
#include <memory>
#include "gtest/gtest.h"
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/module.h"
#include "src/ast/pipeline_stage.h"
#include "src/ast/stage_decoration.h"
#include "src/ast/struct.h"
#include "src/ast/struct_member.h"
#include "src/ast/struct_member_offset_decoration.h"
@ -48,10 +48,12 @@ using MslGeneratorImplTest = testing::Test;
TEST_F(MslGeneratorImplTest, Generate) {
ast::type::VoidType void_type;
ast::Module m;
m.AddFunction(std::make_unique<ast::Function>("my_func", ast::VariableList{},
&void_type));
m.AddEntryPoint(std::make_unique<ast::EntryPoint>(
ast::PipelineStage::kCompute, "", "my_func"));
auto func = std::make_unique<ast::Function>("my_func", ast::VariableList{},
&void_type);
func->add_decoration(
std::make_unique<ast::StageDecoration>(ast::PipelineStage::kCompute));
m.AddFunction(std::move(func));
GeneratorImpl g(&m);

85
src/writer/spirv/builder.cc
View File

@ -185,19 +185,6 @@ bool Builder::Build() {
}
}
// Note, the entry points must be generated after the functions as they need
// to be able to lookup the function information based on the name.
for (const auto& ep : mod_->entry_points()) {
if (!GenerateEntryPoint(ep.get())) {
return false;
}
}
for (const auto& ep : mod_->entry_points()) {
if (!GenerateExecutionModes(ep.get())) {
return false;
}
}
return true;
}
@ -301,52 +288,6 @@ bool Builder::GenerateDiscardStatement(ast::DiscardStatement*) {
return true;
}
bool Builder::GenerateEntryPoint(ast::EntryPoint* ep) {
auto name = ep->name();
if (name.empty()) {
name = ep->function_name();
}
const auto id = id_for_entry_point(ep);
if (id == 0) {
return false;
}
auto stage = pipeline_stage_to_execution_model(ep->stage());
if (stage == SpvExecutionModelMax) {
error_ = "Unknown pipeline stage provided";
return false;
}
OperandList operands = {Operand::Int(stage), Operand::Int(id),
Operand::String(name)};
auto* func = func_name_to_func_[ep->function_name()];
if (func == nullptr) {
error_ = "processing an entry point when the function has not been seen.";
return false;
}
for (const auto* var : func->referenced_module_variables()) {
// For SPIR-V 1.3 we only output Input/output variables. If we update to
// SPIR-V 1.4 or later this should be all variables.
if (var->storage_class() != ast::StorageClass::kInput &&
var->storage_class() != ast::StorageClass::kOutput) {
continue;
}
uint32_t var_id;
if (!scope_stack_.get(var->name(), &var_id)) {
error_ = "unable to find ID for global variable: " + var->name();
return false;
}
operands.push_back(Operand::Int(var_id));
}
push_preamble(spv::Op::OpEntryPoint, operands);
return true;
}
bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
auto stage = pipeline_stage_to_execution_model(func->pipeline_stage());
if (stage == SpvExecutionModelMax) {
@ -378,32 +319,6 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
return true;
}
bool Builder::GenerateExecutionModes(ast::EntryPoint* ep) {
const auto id = id_for_entry_point(ep);
if (id == 0) {
return false;
}
// WGSL fragment shader origin is upper left
if (ep->stage() == ast::PipelineStage::kFragment) {
push_preamble(
spv::Op::OpExecutionMode,
{Operand::Int(id), Operand::Int(SpvExecutionModeOriginUpperLeft)});
} else if (ep->stage() == ast::PipelineStage::kCompute) {
auto* func = func_name_to_func_[ep->function_name()];
uint32_t x = 0;
uint32_t y = 0;
uint32_t z = 0;
std::tie(x, y, z) = func->workgroup_size();
push_preamble(spv::Op::OpExecutionMode,
{Operand::Int(id), Operand::Int(SpvExecutionModeLocalSize),
Operand::Int(x), Operand::Int(y), Operand::Int(z)});
}
return true;
}
bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
// WGSL fragment shader origin is upper left
if (func->pipeline_stage() == ast::PipelineStage::kFragment) {

21
src/writer/spirv/builder.h
View File

@ -185,19 +185,11 @@ class Builder {
/// @returns true if the statement was successfully generated
bool GenerateDiscardStatement(ast::DiscardStatement* stmt);
/// Generates an entry point instruction
/// @param ep the entry point
/// @returns true if the instruction was generated, false otherwise
bool GenerateEntryPoint(ast::EntryPoint* ep);
/// Generates an entry point instruction
/// @param func the function
/// @param id the id of the function
/// @returns true if the instruction was generated, false otherwise
bool GenerateEntryPoint(ast::Function* func, uint32_t id);
/// Generates execution modes for an entry point
/// @param ep the entry point
/// @returns false on failure
bool GenerateExecutionModes(ast::EntryPoint* ep);
/// Generates execution modes for an entry point
/// @param func the function
/// @param id the id of the function
/// @returns false on failure
@ -424,19 +416,6 @@ class Builder {
return func_name_to_id_[name];
}
/// Retrieves the id for an entry point function, or 0 if not found.
/// Emits an error if not found.
/// @param ep the entry point
/// @returns 0 on error
uint32_t id_for_entry_point(ast::EntryPoint* ep) {
auto id = id_for_func_name(ep->function_name());
if (id == 0) {
error_ = "unable to find ID for function: " + ep->function_name();
return 0;
}
return id;
}
ast::Module* mod_;
std::string error_;
uint32_t next_id_ = 1;

242
src/writer/spirv/builder_entry_point_test.cc
View File

@ -1,242 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include <string>
#include "gtest/gtest.h"
#include "spirv/unified1/spirv.h"
#include "spirv/unified1/spirv.hpp11"
#include "src/ast/assignment_statement.h"
#include "src/ast/entry_point.h"
#include "src/ast/function.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/pipeline_stage.h"
#include "src/ast/type/f32_type.h"
#include "src/ast/type/void_type.h"
#include "src/ast/variable.h"
#include "src/ast/workgroup_decoration.h"
#include "src/context.h"
#include "src/type_determiner.h"
#include "src/writer/spirv/builder.h"
#include "src/writer/spirv/spv_dump.h"
namespace tint {
namespace writer {
namespace spirv {
namespace {
// TODO(dsinclair): These have all been ported to stage decorations and this
// whole file can be deleted when we remove EntryPoint.
using BuilderTest = testing::Test;
TEST_F(BuilderTest, EntryPoint) {
ast::type::VoidType void_type;
ast::Function func("frag_main", {}, &void_type);
ast::EntryPoint ep(ast::PipelineStage::kFragment, "main", "frag_main");
ast::Module mod;
Builder b(&mod);
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
ASSERT_TRUE(b.GenerateEntryPoint(&ep)) << b.error();
EXPECT_EQ(DumpInstructions(b.preamble()), R"(OpEntryPoint Fragment %3 "main"
)");
}
TEST_F(BuilderTest, EntryPoint_WithoutName) {
ast::type::VoidType void_type;
ast::Function func("compute_main", {}, &void_type);
ast::EntryPoint ep(ast::PipelineStage::kCompute, "", "compute_main");
ast::Module mod;
Builder b(&mod);
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
ASSERT_TRUE(b.GenerateEntryPoint(&ep)) << b.error();
EXPECT_EQ(DumpInstructions(b.preamble()),
R"(OpEntryPoint GLCompute %3 "compute_main"
)");
}
TEST_F(BuilderTest, EntryPoint_BadFunction) {
ast::EntryPoint ep(ast::PipelineStage::kFragment, "main", "frag_main");
ast::Module mod;
Builder b(&mod);
EXPECT_FALSE(b.GenerateEntryPoint(&ep));
EXPECT_EQ(b.error(), "unable to find ID for function: frag_main");
}
struct EntryPointStageData {
ast::PipelineStage stage;
SpvExecutionModel model;
};
inline std::ostream& operator<<(std::ostream& out, EntryPointStageData data) {
out << data.stage;
return out;
}
using EntryPointStageTest = testing::TestWithParam<EntryPointStageData>;
TEST_P(EntryPointStageTest, Emit) {
auto params = GetParam();
ast::type::VoidType void_type;
ast::Function func("main", {}, &void_type);
ast::EntryPoint ep(params.stage, "", "main");
ast::Module mod;
Builder b(&mod);
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
ASSERT_TRUE(b.GenerateEntryPoint(&ep)) << b.error();
auto preamble = b.preamble();
ASSERT_EQ(preamble.size(), 1u);
EXPECT_EQ(preamble[0].opcode(), spv::Op::OpEntryPoint);
ASSERT_GE(preamble[0].operands().size(), 3u);
EXPECT_EQ(preamble[0].operands()[0].to_i(), params.model);
}
INSTANTIATE_TEST_SUITE_P(
BuilderTest,
EntryPointStageTest,
testing::Values(EntryPointStageData{ast::PipelineStage::kVertex,
SpvExecutionModelVertex},
EntryPointStageData{ast::PipelineStage::kFragment,
SpvExecutionModelFragment},
EntryPointStageData{ast::PipelineStage::kCompute,
SpvExecutionModelGLCompute}));
TEST_F(BuilderTest, EntryPoint_WithUnusedInterfaceIds) {
ast::type::F32Type f32;
ast::type::VoidType void_type;
ast::Function func("main", {}, &void_type);
auto v_in =
std::make_unique<ast::Variable>("my_in", ast::StorageClass::kInput, &f32);
auto v_out = std::make_unique<ast::Variable>(
"my_out", ast::StorageClass::kOutput, &f32);
auto v_wg = std::make_unique<ast::Variable>(
"my_wg", ast::StorageClass::kWorkgroup, &f32);
ast::EntryPoint ep(ast::PipelineStage::kVertex, "", "main");
ast::Module mod;
Builder b(&mod);
EXPECT_TRUE(b.GenerateGlobalVariable(v_in.get())) << b.error();
EXPECT_TRUE(b.GenerateGlobalVariable(v_out.get())) << b.error();
EXPECT_TRUE(b.GenerateGlobalVariable(v_wg.get())) << b.error();
mod.AddGlobalVariable(std::move(v_in));
mod.AddGlobalVariable(std::move(v_out));
mod.AddGlobalVariable(std::move(v_wg));
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
ASSERT_TRUE(b.GenerateEntryPoint(&ep)) << b.error();
EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %1 "my_in"
OpName %4 "my_out"
OpName %7 "my_wg"
OpName %11 "main"
)");
EXPECT_EQ(DumpInstructions(b.types()), R"(%3 = OpTypeFloat 32
%2 = OpTypePointer Input %3
%1 = OpVariable %2 Input
%5 = OpTypePointer Output %3
%6 = OpConstantNull %3
%4 = OpVariable %5 Output %6
%8 = OpTypePointer Workgroup %3
%7 = OpVariable %8 Workgroup
%10 = OpTypeVoid
%9 = OpTypeFunction %10
)");
EXPECT_EQ(DumpInstructions(b.preamble()),
R"(OpEntryPoint Vertex %11 "main"
)");
}
TEST_F(BuilderTest, EntryPoint_WithUsedInterfaceIds) {
ast::type::F32Type f32;
ast::type::VoidType void_type;
ast::Function func("main", {}, &void_type);
auto body = std::make_unique<ast::BlockStatement>();
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("my_out"),
std::make_unique<ast::IdentifierExpression>("my_in")));
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("my_wg"),
std::make_unique<ast::IdentifierExpression>("my_wg")));
// Add duplicate usages so we show they don't get output multiple times.
body->append(std::make_unique<ast::AssignmentStatement>(
std::make_unique<ast::IdentifierExpression>("my_out"),
std::make_unique<ast::IdentifierExpression>("my_in")));
func.set_body(std::move(body));
auto v_in =
std::make_unique<ast::Variable>("my_in", ast::StorageClass::kInput, &f32);
auto v_out = std::make_unique<ast::Variable>(
"my_out", ast::StorageClass::kOutput, &f32);
auto v_wg = std::make_unique<ast::Variable>(
"my_wg", ast::StorageClass::kWorkgroup, &f32);
ast::EntryPoint ep(ast::PipelineStage::kVertex, "", "main");
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
td.RegisterVariableForTesting(v_in.get());
td.RegisterVariableForTesting(v_out.get());
td.RegisterVariableForTesting(v_wg.get());
ASSERT_TRUE(td.DetermineFunction(&func)) << td.error();
Builder b(&mod);
EXPECT_TRUE(b.GenerateGlobalVariable(v_in.get())) << b.error();
EXPECT_TRUE(b.GenerateGlobalVariable(v_out.get())) << b.error();
EXPECT_TRUE(b.GenerateGlobalVariable(v_wg.get())) << b.error();
mod.AddGlobalVariable(std::move(v_in));
mod.AddGlobalVariable(std::move(v_out));
mod.AddGlobalVariable(std::move(v_wg));
ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
ASSERT_TRUE(b.GenerateEntryPoint(&ep)) << b.error();
EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %1 "my_in"
OpName %4 "my_out"
OpName %7 "my_wg"
OpName %11 "main"
)");
EXPECT_EQ(DumpInstructions(b.types()), R"(%3 = OpTypeFloat 32
%2 = OpTypePointer Input %3
%1 = OpVariable %2 Input
%5 = OpTypePointer Output %3
%6 = OpConstantNull %3
%4 = OpVariable %5 Output %6
%8 = OpTypePointer Workgroup %3
%7 = OpVariable %8 Workgroup
%10 = OpTypeVoid
%9 = OpTypeFunction %10
)");
EXPECT_EQ(DumpInstructions(b.preamble()),
R"(OpEntryPoint Vertex %11 "main" %4 %1
)");
}
} // namespace
} // namespace spirv
} // namespace writer
} // namespace tint

48
src/writer/wgsl/generator_impl.cc
View File

@ -84,14 +84,6 @@ bool GeneratorImpl::Generate(const ast::Module& module) {
out_ << std::endl;
}
for (const auto& ep : module.entry_points()) {
if (!EmitEntryPoint(ep.get())) {
return false;
}
}
if (!module.entry_points().empty())
out_ << std::endl;
for (auto* const alias : module.alias_types()) {
if (!EmitAliasType(alias)) {
return false;
@ -132,27 +124,8 @@ bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
out_ << std::endl;
}
bool found_entry_point = false;
std::string ep_function_name = "";
for (const auto& ep : module.entry_points()) {
std::string ep_name = ep->name();
if (ep_name.empty()) {
ep_name = ep->function_name();
}
ep_function_name = ep->function_name();
if (ep->stage() != stage || ep_name != name) {
continue;
}
if (!EmitEntryPoint(ep.get())) {
return false;
}
found_entry_point = true;
break;
}
out_ << std::endl;
if (!found_entry_point) {
auto* func = module.FindFunctionByNameAndStage(name, stage);
if (func == nullptr) {
error_ = "Unable to find requested entry point: " + name;
return false;
}
@ -178,12 +151,6 @@ bool GeneratorImpl::GenerateEntryPoint(const ast::Module& module,
}
}
auto* func = module.FindFunctionByName(ep_function_name);
if (!func) {
error_ = "Unable to find entry point function: " + ep_function_name;
return false;
}
bool found_func_variable = false;
for (auto* var : func->referenced_module_variables()) {
if (!EmitVariable(var)) {
@ -225,17 +192,6 @@ bool GeneratorImpl::EmitAliasType(const ast::type::AliasType* alias) {
return true;
}
bool GeneratorImpl::EmitEntryPoint(const ast::EntryPoint* ep) {
make_indent();
out_ << "entry_point " << ep->stage() << " ";
if (!ep->name().empty() && ep->name() != ep->function_name()) {
out_ << R"(as ")" << ep->name() << R"(" )";
}
out_ << "= " << ep->function_name() << ";" << std::endl;
return true;
}
bool GeneratorImpl::EmitExpression(ast::Expression* expr) {
if (expr->IsArrayAccessor()) {
return EmitArrayAccessor(expr->AsArrayAccessor());

5
src/writer/wgsl/generator_impl.h
View File

@ -20,7 +20,6 @@
#include "src/ast/array_accessor_expression.h"
#include "src/ast/constructor_expression.h"
#include "src/ast/entry_point.h"
#include "src/ast/identifier_expression.h"
#include "src/ast/import.h"
#include "src/ast/module.h"
@ -116,10 +115,6 @@ class GeneratorImpl : public TextGenerator {
/// @param stmt the statement to emit
/// @returns true if the statement was emitted
bool EmitElse(ast::ElseStatement* stmt);
/// Handles generating an entry_point command
/// @param ep the entry point
/// @returns true if the entry point was emitted
bool EmitEntryPoint(const ast::EntryPoint* ep);
/// Handles generate an Expression
/// @param expr the expression
/// @returns true if the expression was emitted

46
src/writer/wgsl/generator_impl_entry_point_test.cc
View File

@ -1,46 +0,0 @@
// Copyright 2020 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "gtest/gtest.h"
#include "src/writer/wgsl/generator_impl.h"
namespace tint {
namespace writer {
namespace wgsl {
namespace {
using WgslGeneratorImplTest = testing::Test;
TEST_F(WgslGeneratorImplTest, EmitEntryPoint_NoName) {
ast::EntryPoint ep(ast::PipelineStage::kFragment, "", "frag_main");
GeneratorImpl g;
ASSERT_TRUE(g.EmitEntryPoint(&ep)) << g.error();
EXPECT_EQ(g.result(), R"(entry_point fragment = frag_main;
)");
}
TEST_F(WgslGeneratorImplTest, EmitEntryPoint_WithName) {
ast::EntryPoint ep(ast::PipelineStage::kFragment, "main", "frag_main");
GeneratorImpl g;
ASSERT_TRUE(g.EmitEntryPoint(&ep)) << g.error();
EXPECT_EQ(g.result(), R"(entry_point fragment as "main" = frag_main;
)");
}
} // namespace
} // namespace wgsl
} // namespace writer
} // namespace tint

11
test/compute_boids.wgsl
View File

@ -21,7 +21,8 @@ import "GLSL.std.450" as std;
[[location(2)]] var<in> a_pos : vec2<f32>;
[[builtin(position)]] var<out> gl_Position : vec4<f32>;
fn vtx_main() -> void {
[[stage(vertex)]]
fn vert_main() -> void {
var angle : f32 = -std::atan2(a_particleVel.x, a_particleVel.y);
var pos : vec2<f32> = vec2<f32>(
(a_pos.x * std::cos(angle)) - (a_pos.y * std::sin(angle)),
@ -29,16 +30,15 @@ fn vtx_main() -> void {
gl_Position = vec4<f32>(pos + a_particlePos, 0.0, 1.0);
return;
}
entry_point vertex as "vert_main" = vtx_main;
# fragment shader
[[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]]
fn frag_main() -> void {
fragColor = vec4<f32>(1.0, 1.0, 1.0, 1.0);
return;
}
entry_point fragment as "frag_main" = frag_main;
# compute shader
type Particle = [[block]] struct {
@ -67,7 +67,8 @@ type Particles = [[block]] struct {
[[builtin(global_invocation_id)]] var<in> gl_GlobalInvocationID : vec3<u32>;
# https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
fn compute_main() -> void {
[[stage(compute)]]
fn comp_main() -> void {
var index : u32 = gl_GlobalInvocationID.x;
if (index >= 5) {
return;
@ -149,5 +150,3 @@ fn compute_main() -> void {
return;
}
entry_point compute as "comp_main" = compute_main;

5
test/cube.wgsl
View File

@ -12,9 +12,6 @@
# See the License for the specific language governing permissions and
# limitations under the License.
entry_point vertex = vtx_main;
entry_point fragment = frag_main;
# Vertex shader
type Uniforms = [[block]] struct {
[[offset(0)]] modelViewProjectionMatrix : mat4x4<f32>;
@ -27,6 +24,7 @@ type Uniforms = [[block]] struct {
[[location(0)]] var<out> vtxFragColor : vec4<f32>;
[[builtin(position)]] var<out> Position : vec4<f32>;
[[stage(vertex)]]
fn vtx_main() -> void {
Position = uniforms.modelViewProjectionMatrix * cur_position;
vtxFragColor = color;
@ -37,6 +35,7 @@ fn vtx_main() -> void {
[[location(0)]] var<in> fragColor : vec4<f32>;
[[location(0)]] var<out> outColor : vec4<f32>;
[[stage(fragment)]]
fn frag_main() -> void {
outColor = fragColor;
return;

2
test/simple.wgsl
View File

@ -18,10 +18,10 @@ fn bar() -> void {
return;
}
[[stage(fragment)]]
fn main() -> void {
var a : vec2<f32> = vec2<f32>();
gl_FragColor = vec4<f32>(0.4, 0.4, 0.8, 1.0);
bar();
return;
}
entry_point fragment = main;

2
test/switch-case-selector-must-have-the-same-type-as-the-selector-expression-2.fail.wgsl
View File

@ -1,6 +1,6 @@
# v-switch03: line 7: the case selector values must have the same type as the selector expression
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: i32 = -2;
switch (a) {

2
test/switch-case-selector-must-have-the-same-type-as-the-selector-expression.fail.wgsl
View File

@ -1,6 +1,6 @@
# v-switch03: line 7: the case selector values must have the same type as the selector expression
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: u32 = 2;
switch (a) {

2
test/switch-case-selector-value-must-be-unique.fail.wgsl
View File

@ -1,7 +1,7 @@
# v-switch04: line 9: a literal value must not appear more than once in the case selectors for a
# switch statement: '0'
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: u32 = 2;
switch (a) {

2
test/switch-fallthrough-must-not-be-last-stmt-of-last-clause.fail.wgsl
View File

@ -1,7 +1,7 @@
# v-switch05: line 9: a fallthrough statement must not appear as the last statement in last clause
# of a switch
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: i32 = -2;
switch (a) {

2
test/switch-must-have-exactly-one-default-clause-2.fail.wgsl
View File

@ -1,6 +1,6 @@
# v-0008: line 6: switch statement must have exactly one default clause
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: i32 = 2;
switch (a) {

2
test/switch-must-have-exactly-one-default-clause.fail.wgsl
View File

@ -1,6 +1,6 @@
# v-0008: line 6: switch statement must have exactly one default clause
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: i32 = 2;
switch (a) {

2
test/switch-selector-expression-must-be-scalar-integer-type.fail.wgsl
View File

@ -1,6 +1,6 @@
# v-switch01: line 6: switch statement selector expression must be of a scalar integer type
entry_point vertex = main;
[[stage(vertex)]]
fn main() -> void {
var a: f32 = 3.14;
switch (a) {

5
test/triangle.wgsl
View File

@ -21,16 +21,17 @@ const pos : array<vec2<f32>, 3> = array<vec2<f32>, 3>(
[[builtin(position)]] var<out> Position : vec4<f32>;
[[builtin(vertex_idx)]] var<in> VertexIndex : i32;
[[stage(vertex)]]
fn vtx_main() -> void {
Position = vec4<f32>(pos[VertexIndex], 0.0, 1.0);
return;
}
entry_point vertex as "main" = vtx_main;
# Fragment shader
[[location(0)]] var<out> outColor : vec4<f32>;
[[stage(fragment)]]
fn frag_main() -> void {
outColor = vec4<f32>(1.0, 0.0, 0.0, 1.0);
return;
}
entry_point fragment = frag_main;

2
test/undefined-variable-global-scope.wgsl
View File

@ -15,9 +15,9 @@
# Vertex shader
var<private> a : i32;
[[stage(vertex)]]
fn main() -> void {
a = 2;
return;
}
entry_point vertex as "main" = main;

2
test/undefined-variable-inner-block.wgsl
View File

@ -12,9 +12,9 @@
# See the License for the specific language governing permissions and
# limitations under the License.
[[stage(fragment)]]
fn main() -> void {
var a : f32 = 2.0;
{ a = 3.14;}
return;
}
entry_point fragment = main;