[msl-writer] Add import emission.

This CL adds emission of import function calls to the MSL backend.

Bug: tint:8
Change-Id: Ib8b8638b11caee2ff3557d551447b215ef2a4c69
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/25262
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
dan sinclair 2020-07-20 22:13:00 +00:00
parent 747e80a9b3
commit 26c8a216e8
5 changed files with 491 additions and 6 deletions

View File

@ -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",

View File

@ -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

View File

@ -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();

View File

@ -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

View File

@ -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 <memory>
#include <string>
#include <vector>
#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<MslImportData>;
TEST_P(MslImportData_SingleParamTest, FloatScalar) {
auto param = GetParam();
ast::type::F32Type f32;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 1)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "sabs"}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<MslImportData>;
TEST_P(MslImportData_DualParamTest, FloatScalar) {
auto param = GetParam();
ast::type::F32Type f32;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 2.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<MslImportData>;
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<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
type_params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 2.f)));
type_params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 3.f)));
ast::ExpressionList params;
params.push_back(std::make_unique<ast::TypeConstructorExpression>(
&vec, std::move(type_params)));
type_params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 4.f)));
type_params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 5.f)));
type_params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 6.f)));
params.push_back(std::make_unique<ast::TypeConstructorExpression>(
&vec, std::move(type_params)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<MslImportData>;
TEST_P(MslImportData_DualParam_Int_Test, IntScalar) {
auto param = GetParam();
ast::type::I32Type i32;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 1)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 2)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<MslImportData>;
TEST_P(MslImportData_TripleParamTest, FloatScalar) {
auto param = GetParam();
ast::type::F32Type f32;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 1.f)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 2.f)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::FloatLiteral>(&f32, 3.f)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<MslImportData>;
TEST_P(MslImportData_TripleParam_Int_Test, IntScalar) {
auto param = GetParam();
ast::type::I32Type i32;
ast::ExpressionList params;
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 1)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 2)));
params.push_back(std::make_unique<ast::ScalarConstructorExpression>(
std::make_unique<ast::SintLiteral>(&i32, 3)));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", param.name}),
std::move(params));
Context ctx;
ast::Module mod;
TypeDeterminer td(&ctx, &mod);
mod.AddImport(std::make_unique<ast::Import>("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<ast::Variable>(
"var", ast::StorageClass::kFunction, &mat);
ast::ExpressionList params;
params.push_back(std::make_unique<ast::IdentifierExpression>("var"));
ast::CallExpression expr(std::make_unique<ast::IdentifierExpression>(
std::vector<std::string>{"std", "determinant"}),
std::move(params));
Context ctx;
ast::Module mod;
mod.AddGlobalVariable(std::move(var));
mod.AddImport(std::make_unique<ast::Import>("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