diff --git a/BUILD.gn b/BUILD.gn index a1a08410fb..9bd4307655 100644 --- a/BUILD.gn +++ b/BUILD.gn @@ -949,6 +949,7 @@ source_set("tint_unittests_msl_writer_src") { "src/writer/msl/generator_impl_function_test.cc", "src/writer/msl/generator_impl_identifier_test.cc", "src/writer/msl/generator_impl_if_test.cc", + "src/writer/msl/generator_impl_import_test.cc", "src/writer/msl/generator_impl_kill_test.cc", "src/writer/msl/generator_impl_loop_test.cc", "src/writer/msl/generator_impl_member_accessor_test.cc", diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 43033f904e..fe9486d039 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -509,6 +509,7 @@ if(${TINT_BUILD_MSL_WRITER}) writer/msl/generator_impl_function_test.cc writer/msl/generator_impl_identifier_test.cc writer/msl/generator_impl_if_test.cc + writer/msl/generator_impl_import_test.cc writer/msl/generator_impl_kill_test.cc writer/msl/generator_impl_loop_test.cc writer/msl/generator_impl_member_accessor_test.cc diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 2c6411e462..e9f85d55bc 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -14,6 +14,7 @@ #include "src/writer/msl/generator_impl.h" +#include "spirv/unified1/GLSL.std.450.h" #include "src/ast/array_accessor_expression.h" #include "src/ast/as_expression.h" #include "src/ast/assignment_statement.h" @@ -477,13 +478,130 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { out_ << ")"; } else { - // TODO(dsinclair): Handle imported function - error_ = "imported calls not handled yet"; - return false; + return EmitImportFunction(expr); } return true; } +bool GeneratorImpl::EmitImportFunction(ast::CallExpression* expr) { + auto* ident = expr->func()->AsIdentifier(); + + auto* imp = module_->FindImportByName(ident->path()); + if (imp == nullptr) { + error_ = "unable to find import for " + ident->path(); + return 0; + } + auto id = imp->GetIdForMethod(ident->name()); + if (id == 0) { + error_ = "unable to lookup: " + ident->name() + " in " + ident->path(); + } + + out_ << "metal::"; + switch (id) { + case GLSLstd450Acos: + case GLSLstd450Acosh: + case GLSLstd450Asin: + case GLSLstd450Asinh: + case GLSLstd450Atan: + case GLSLstd450Atan2: + case GLSLstd450Atanh: + case GLSLstd450Ceil: + case GLSLstd450Cos: + case GLSLstd450Cosh: + case GLSLstd450Cross: + case GLSLstd450Determinant: + case GLSLstd450Distance: + case GLSLstd450Exp: + case GLSLstd450Exp2: + case GLSLstd450FAbs: + case GLSLstd450FaceForward: + case GLSLstd450Floor: + case GLSLstd450Fma: + case GLSLstd450Fract: + case GLSLstd450Length: + case GLSLstd450Log: + case GLSLstd450Log2: + case GLSLstd450Normalize: + case GLSLstd450Pow: + case GLSLstd450Reflect: + case GLSLstd450Round: + case GLSLstd450Sin: + case GLSLstd450Sinh: + case GLSLstd450SmoothStep: + case GLSLstd450Sqrt: + case GLSLstd450Step: + case GLSLstd450Tan: + case GLSLstd450Tanh: + case GLSLstd450Trunc: + out_ << ident->name(); + break; + case GLSLstd450InverseSqrt: + out_ << "rsqrt"; + break; + case GLSLstd450FMax: + case GLSLstd450NMax: + out_ << "fmax"; + break; + case GLSLstd450FMin: + case GLSLstd450NMin: + out_ << "fmin"; + break; + case GLSLstd450FMix: + out_ << "mix"; + break; + case GLSLstd450FSign: + out_ << "sign"; + break; + case GLSLstd450SAbs: + out_ << "abs"; + break; + case GLSLstd450SMax: + case GLSLstd450UMax: + out_ << "max"; + break; + case GLSLstd450SMin: + case GLSLstd450UMin: + out_ << "min"; + break; + case GLSLstd450FClamp: + case GLSLstd450SClamp: + case GLSLstd450NClamp: + case GLSLstd450UClamp: + out_ << "clamp"; + break; + // TODO(dsinclair): Determine mappings for the following + case GLSLstd450Degrees: + case GLSLstd450FindILsb: + case GLSLstd450FindUMsb: + case GLSLstd450FindSMsb: + case GLSLstd450InterpolateAtCentroid: + case GLSLstd450MatrixInverse: + case GLSLstd450Radians: + case GLSLstd450RoundEven: + case GLSLstd450SSign: + default: + error_ = "Unknown import method: " + ident->name(); + return false; + } + + out_ << "("; + bool first = true; + const auto& params = expr->params(); + for (const auto& param : params) { + if (!first) { + out_ << ", "; + } + first = false; + + if (!EmitExpression(param.get())) { + return false; + } + } + out_ << ")"; + + return true; +} + bool GeneratorImpl::EmitCase(ast::CaseStatement* stmt) { make_indent(); diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h index 5578a6226f..cc7d33ad1a 100644 --- a/src/writer/msl/generator_impl.h +++ b/src/writer/msl/generator_impl.h @@ -83,7 +83,7 @@ class GeneratorImpl : public TextGenerator { bool EmitCall(ast::CallExpression* expr); /// Handles a case statement /// @param stmt the statement - /// @returns true if the statment was emitted successfully + /// @returns true if the statement was emitted successfully bool EmitCase(ast::CaseStatement* stmt); /// Handles generating a cast expression /// @param expr the cast expression @@ -128,12 +128,16 @@ class GeneratorImpl : public TextGenerator { const std::string& ep_name); /// Handles generating an identifier expression /// @param expr the identifier expression - /// @returns true if the identifeir was emitted + /// @returns true if the identifier was emitted bool EmitIdentifier(ast::IdentifierExpression* expr); /// Handles an if statement /// @param stmt the statement to emit /// @returns true if the statement was successfully emitted bool EmitIf(ast::IfStatement* stmt); + /// Handles genreating an import expression + /// @param expr the expression + /// @returns true if the expression was successfully emitted. + bool EmitImportFunction(ast::CallExpression* expr); /// Handles generating a kill statement /// @param stmt the kill statement /// @returns true if the statement was successfully emitted @@ -144,7 +148,7 @@ class GeneratorImpl : public TextGenerator { bool EmitLiteral(ast::Literal* lit); /// Handles a loop statement /// @param stmt the statement to emit - /// @returns true if the statement was emtited + /// @returns true if the statement was emitted bool EmitLoop(ast::LoopStatement* stmt); /// Handles a member accessor expression /// @param expr the member accessor expression diff --git a/src/writer/msl/generator_impl_import_test.cc b/src/writer/msl/generator_impl_import_test.cc new file mode 100644 index 0000000000..fcf0d50b02 --- /dev/null +++ b/src/writer/msl/generator_impl_import_test.cc @@ -0,0 +1,361 @@ +// 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 +#include +#include + +#include "gtest/gtest.h" +#include "src/ast/call_expression.h" +#include "src/ast/float_literal.h" +#include "src/ast/identifier_expression.h" +#include "src/ast/module.h" +#include "src/ast/scalar_constructor_expression.h" +#include "src/ast/sint_literal.h" +#include "src/ast/type/f32_type.h" +#include "src/ast/type/i32_type.h" +#include "src/ast/type/matrix_type.h" +#include "src/ast/type/vector_type.h" +#include "src/ast/type_constructor_expression.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; + +struct MslImportData { + const char* name; + const char* msl_name; +}; +inline std::ostream& operator<<(std::ostream& out, MslImportData data) { + out << data.name; + return out; +} +using MslImportData_SingleParamTest = testing::TestWithParam; +TEST_P(MslImportData_SingleParamTest, FloatScalar) { + auto param = GetParam(); + + ast::type::F32Type f32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&f32, 1.f))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), + std::string("metal::") + param.msl_name + "(1.00000000f)"); +} +INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, + MslImportData_SingleParamTest, + testing::Values(MslImportData{"acos", "acos"}, + MslImportData{"acosh", "acosh"}, + MslImportData{"asin", "asin"}, + MslImportData{"asinh", "asinh"}, + MslImportData{"atan", "atan"}, + MslImportData{"atanh", "atanh"}, + MslImportData{"cos", "cos"}, + MslImportData{"cosh", "cosh"}, + MslImportData{"ceil", "ceil"}, + MslImportData{"exp", "exp"}, + MslImportData{"exp2", "exp2"}, + MslImportData{"fabs", "fabs"}, + MslImportData{"floor", "floor"}, + MslImportData{"fract", "fract"}, + MslImportData{"inversesqrt", "rsqrt"}, + MslImportData{"length", "length"}, + MslImportData{"log", "log"}, + MslImportData{"log2", "log2"}, + MslImportData{"normalize", + "normalize"}, + MslImportData{"round", "round"}, + MslImportData{"fsign", "sign"}, + MslImportData{"sin", "sin"}, + MslImportData{"sinh", "sinh"}, + MslImportData{"sqrt", "sqrt"}, + MslImportData{"tan", "tan"}, + MslImportData{"tanh", "tanh"}, + MslImportData{"trunc", "trunc"})); + +TEST_F(MslGeneratorImplTest, MslImportData_SingleParamTest_IntScalar) { + ast::type::I32Type i32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&i32, 1))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", "sabs"}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), R"(metal::abs(1))"); +} + +using MslImportData_DualParamTest = testing::TestWithParam; +TEST_P(MslImportData_DualParamTest, FloatScalar) { + auto param = GetParam(); + + ast::type::F32Type f32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&f32, 1.f))); + params.push_back(std::make_unique( + std::make_unique(&f32, 2.f))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + + "(1.00000000f, 2.00000000f)"); +} +INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, + MslImportData_DualParamTest, + testing::Values(MslImportData{"atan2", "atan2"}, + MslImportData{"distance", "distance"}, + MslImportData{"fmax", "fmax"}, + MslImportData{"fmin", "fmin"}, + MslImportData{"nmax", "fmax"}, + MslImportData{"nmin", "fmin"}, + MslImportData{"pow", "pow"}, + MslImportData{"reflect", "reflect"}, + MslImportData{"step", "step"})); + +using MslImportData_DualParam_VectorTest = + testing::TestWithParam; +TEST_P(MslImportData_DualParam_VectorTest, FloatVector) { + auto param = GetParam(); + + ast::type::F32Type f32; + ast::type::VectorType vec(&f32, 3); + + ast::ExpressionList type_params; + type_params.push_back(std::make_unique( + std::make_unique(&f32, 1.f))); + type_params.push_back(std::make_unique( + std::make_unique(&f32, 2.f))); + type_params.push_back(std::make_unique( + std::make_unique(&f32, 3.f))); + + ast::ExpressionList params; + params.push_back(std::make_unique( + &vec, std::move(type_params))); + + type_params.push_back(std::make_unique( + std::make_unique(&f32, 4.f))); + type_params.push_back(std::make_unique( + std::make_unique(&f32, 5.f))); + type_params.push_back(std::make_unique( + std::make_unique(&f32, 6.f))); + params.push_back(std::make_unique( + &vec, std::move(type_params))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + + "(float3(1.00000000f, 2.00000000f, 3.00000000f), " + "float3(4.00000000f, 5.00000000f, 6.00000000f))"); +} +INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, + MslImportData_DualParam_VectorTest, + testing::Values(MslImportData{"cross", "cross"})); + +using MslImportData_DualParam_Int_Test = testing::TestWithParam; +TEST_P(MslImportData_DualParam_Int_Test, IntScalar) { + auto param = GetParam(); + + ast::type::I32Type i32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&i32, 1))); + params.push_back(std::make_unique( + std::make_unique(&i32, 2))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + "(1, 2)"); +} +INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, + MslImportData_DualParam_Int_Test, + testing::Values(MslImportData{"smax", "max"}, + MslImportData{"smin", "min"}, + MslImportData{"umax", "max"}, + MslImportData{"umin", "min"})); + +using MslImportData_TripleParamTest = testing::TestWithParam; +TEST_P(MslImportData_TripleParamTest, FloatScalar) { + auto param = GetParam(); + + ast::type::F32Type f32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&f32, 1.f))); + params.push_back(std::make_unique( + std::make_unique(&f32, 2.f))); + params.push_back(std::make_unique( + std::make_unique(&f32, 3.f))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + + "(1.00000000f, 2.00000000f, 3.00000000f)"); +} +INSTANTIATE_TEST_SUITE_P( + MslGeneratorImplTest, + MslImportData_TripleParamTest, + testing::Values(MslImportData{"faceforward", "faceforward"}, + MslImportData{"fma", "fma"}, + MslImportData{"fclamp", "clamp"}, + MslImportData{"fmix", "mix"}, + MslImportData{"nclamp", "clamp"}, + MslImportData{"smoothstep", "smoothstep"})); + +using MslImportData_TripleParam_Int_Test = + testing::TestWithParam; +TEST_P(MslImportData_TripleParam_Int_Test, IntScalar) { + auto param = GetParam(); + + ast::type::I32Type i32; + + ast::ExpressionList params; + params.push_back(std::make_unique( + std::make_unique(&i32, 1))); + params.push_back(std::make_unique( + std::make_unique(&i32, 2))); + params.push_back(std::make_unique( + std::make_unique(&i32, 3))); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", param.name}), + std::move(params)); + + Context ctx; + ast::Module mod; + TypeDeterminer td(&ctx, &mod); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::") + param.msl_name + "(1, 2, 3)"); +} +INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, + MslImportData_TripleParam_Int_Test, + testing::Values(MslImportData{"sclamp", "clamp"}, + MslImportData{"uclamp", "clamp"})); + +TEST_F(MslGeneratorImplTest, MslImportData_Determinant) { + ast::type::F32Type f32; + ast::type::MatrixType mat(&f32, 3, 3); + + auto var = std::make_unique( + "var", ast::StorageClass::kFunction, &mat); + + ast::ExpressionList params; + params.push_back(std::make_unique("var")); + + ast::CallExpression expr(std::make_unique( + std::vector{"std", "determinant"}), + std::move(params)); + + Context ctx; + ast::Module mod; + mod.AddGlobalVariable(std::move(var)); + mod.AddImport(std::make_unique("GLSL.std.450", "std")); + + TypeDeterminer td(&ctx, &mod); + // Register the global + ASSERT_TRUE(td.Determine()) << td.error(); + ASSERT_TRUE(td.DetermineResultType(&expr)) << td.error(); + + GeneratorImpl g(&mod); + ASSERT_TRUE(g.EmitImportFunction(&expr)) << g.error(); + EXPECT_EQ(g.result(), std::string("metal::determinant(var)")); +} + +} // namespace +} // namespace msl +} // namespace writer +} // namespace tint