From 2aaf7b5c766fd7e35074401979946fcfbde70321 Mon Sep 17 00:00:00 2001 From: dan sinclair <dsinclair@chromium.org> Date: Thu, 25 Feb 2021 21:17:47 +0000 Subject: [PATCH] [MslWriter] Emit using statement for metal This CL updates the MSL backend to emit a using namespace for metal instead of using the `metal::` prefix. Bug: tint:463 Change-Id: I63d3ea5b5a56e61d71cd6d17a51a5120363ea007 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/42741 Auto-Submit: dan sinclair <dsinclair@chromium.org> Reviewed-by: Ryan Harrison <rharrison@chromium.org> Reviewed-by: Ben Clayton <bclayton@google.com> Commit-Queue: dan sinclair <dsinclair@chromium.org> --- src/writer/msl/generator_impl.cc | 3 +- .../msl/generator_impl_function_test.cc | 20 +++ src/writer/msl/generator_impl_import_test.cc | 20 +-- .../msl/generator_impl_intrinsic_test.cc | 150 ++++++++---------- src/writer/msl/generator_impl_test.cc | 1 + 5 files changed, 101 insertions(+), 93 deletions(-) diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index a603d07688..dee90cbeaa 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -122,6 +122,7 @@ std::string GeneratorImpl::generate_name(const std::string& prefix) { bool GeneratorImpl::Generate() { out_ << "#include <metal_stdlib>" << std::endl << std::endl; + out_ << "using namespace metal;" << std::endl; for (auto* global : program_->AST().GlobalVariables()) { auto* sem = program_->Sem().Get(global); @@ -789,7 +790,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, std::string GeneratorImpl::generate_builtin_name( const semantic::Intrinsic* intrinsic) { - std::string out = "metal::"; + std::string out = ""; switch (intrinsic->Type()) { case semantic::IntrinsicType::kAcos: case semantic::IntrinsicType::kAll: diff --git a/src/writer/msl/generator_impl_function_test.cc b/src/writer/msl/generator_impl_function_test.cc index 69f2ced1e9..714c0b20bf 100644 --- a/src/writer/msl/generator_impl_function_test.cc +++ b/src/writer/msl/generator_impl_function_test.cc @@ -69,6 +69,7 @@ TEST_F(MslGeneratorImplTest, Emit_Function) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; void my_func() { return; } @@ -90,6 +91,7 @@ TEST_F(MslGeneratorImplTest, Emit_Function_Name_Collision) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; void main_tint_0() { return; } @@ -115,6 +117,7 @@ TEST_F(MslGeneratorImplTest, Emit_Function_WithParams) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; void my_func(float a, int b) { return; } @@ -133,6 +136,7 @@ TEST_F(MslGeneratorImplTest, Emit_FunctionDecoration_EntryPoint_NoReturn_Void) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; fragment void main_tint_0() { return; } @@ -160,6 +164,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct main_in { float foo [[user(locn0)]]; }; @@ -197,6 +202,7 @@ TEST_F(MslGeneratorImplTest, Emit_FunctionDecoration_EntryPoint_WithInOutVars) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct frag_main_in { float foo [[user(locn0)]]; }; @@ -240,6 +246,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct frag_main_out { float depth [[depth(any)]]; }; @@ -275,6 +282,7 @@ TEST_F(MslGeneratorImplTest, Emit_FunctionDecoration_EntryPoint_With_Uniform) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; fragment void frag_main(constant float4& coord [[buffer(0)]]) { float v = coord.x; return; @@ -316,6 +324,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct Data { int a; float b; @@ -361,6 +370,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct Data { int a; float b; @@ -411,6 +421,7 @@ TEST_F( ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct ep_1_in { float foo [[user(locn0)]]; }; @@ -465,6 +476,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct ep_1_out { float depth [[depth(any)]]; }; @@ -519,6 +531,7 @@ TEST_F( ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct ep_1_out { float depth [[depth(any)]]; }; @@ -572,6 +585,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; float sub_func(constant float4& coord, float param) { return coord.x; } @@ -624,6 +638,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct Data { int a; float b; @@ -684,6 +699,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct Data { int a; float b; @@ -728,6 +744,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct ep_1_out { float bar [[color(1)]]; }; @@ -756,6 +773,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; kernel void main_tint_0() { return; } @@ -780,6 +798,7 @@ TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; void my_func(float a[5]) { return; } @@ -850,6 +869,7 @@ TEST_F(MslGeneratorImplTest, ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; struct Data { float d; }; diff --git a/src/writer/msl/generator_impl_import_test.cc b/src/writer/msl/generator_impl_import_test.cc index 4894db3ed4..5176a92b8d 100644 --- a/src/writer/msl/generator_impl_import_test.cc +++ b/src/writer/msl/generator_impl_import_test.cc @@ -65,8 +65,7 @@ TEST_P(MslImportData_SingleParamTest, FloatScalar) { auto* intrinsic = target->As<semantic::Intrinsic>(); ASSERT_NE(intrinsic, nullptr); - ASSERT_EQ(gen.generate_builtin_name(intrinsic), - std::string("metal::") + param.msl_name); + ASSERT_EQ(gen.generate_builtin_name(intrinsic), param.msl_name); } INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, MslImportData_SingleParamTest, @@ -101,7 +100,7 @@ TEST_F(MslGeneratorImplTest, MslImportData_SingleParamTest_IntScalar) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), R"(metal::abs(1))"); + EXPECT_EQ(gen.result(), R"(abs(1))"); } using MslImportData_DualParamTest = TestParamHelper<MslImportData>; @@ -114,8 +113,7 @@ TEST_P(MslImportData_DualParamTest, FloatScalar) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), - std::string("metal::") + param.msl_name + "(1.0f, 2.0f)"); + EXPECT_EQ(gen.result(), std::string(param.msl_name) + "(1.0f, 2.0f)"); } INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, MslImportData_DualParamTest, @@ -138,7 +136,7 @@ TEST_P(MslImportData_DualParam_VectorTest, FloatVector) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), std::string("metal::") + param.msl_name + + EXPECT_EQ(gen.result(), std::string(param.msl_name) + "(float3(1.0f, 2.0f, 3.0f), " "float3(4.0f, 5.0f, 6.0f))"); } @@ -156,7 +154,7 @@ TEST_P(MslImportData_DualParam_Int_Test, IntScalar) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), std::string("metal::") + param.msl_name + "(1, 2)"); + EXPECT_EQ(gen.result(), std::string(param.msl_name) + "(1, 2)"); } INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, MslImportData_DualParam_Int_Test, @@ -173,8 +171,7 @@ TEST_P(MslImportData_TripleParamTest, FloatScalar) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), - std::string("metal::") + param.msl_name + "(1.0f, 2.0f, 3.0f)"); + EXPECT_EQ(gen.result(), std::string(param.msl_name) + "(1.0f, 2.0f, 3.0f)"); } INSTANTIATE_TEST_SUITE_P( MslGeneratorImplTest, @@ -195,8 +192,7 @@ TEST_P(MslImportData_TripleParam_Int_Test, IntScalar) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), - std::string("metal::") + param.msl_name + "(1, 2, 3)"); + EXPECT_EQ(gen.result(), std::string(param.msl_name) + "(1, 2, 3)"); } INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest, MslImportData_TripleParam_Int_Test, @@ -213,7 +209,7 @@ TEST_F(MslGeneratorImplTest, MslImportData_Determinant) { GeneratorImpl& gen = Build(); ASSERT_TRUE(gen.EmitCall(expr)) << gen.error(); - EXPECT_EQ(gen.result(), std::string("metal::determinant(var)")); + EXPECT_EQ(gen.result(), std::string("determinant(var)")); } } // namespace diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc index cd28dce7b5..7f263636a1 100644 --- a/src/writer/msl/generator_impl_intrinsic_test.cc +++ b/src/writer/msl/generator_impl_intrinsic_test.cc @@ -199,98 +199,88 @@ INSTANTIATE_TEST_SUITE_P( MslGeneratorImplTest, MslIntrinsicTest, testing::Values( - IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "metal::fabs"}, - IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "metal::abs"}, - IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "metal::acos"}, - IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "metal::all"}, - IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "metal::any"}, - IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "metal::asin"}, - IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "metal::atan"}, - IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "metal::atan2"}, - IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "metal::ceil"}, - IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "metal::clamp"}, - IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "metal::clamp"}, - IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "metal::cos"}, - IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "metal::cosh"}, + IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "fabs"}, + IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "abs"}, + IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "acos"}, + IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "all"}, + IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "any"}, + IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "asin"}, + IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "atan"}, + IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "atan2"}, + IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "ceil"}, + IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "clamp"}, + IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "clamp"}, + IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "cos"}, + IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "cosh"}, IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32, - "metal::popcount"}, - IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "metal::cross"}, + "popcount"}, + IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "cross"}, IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32, - "metal::determinant"}, - IntrinsicData{IntrinsicType::kDistance, ParamType::kF32, - "metal::distance"}, - IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "metal::dot"}, - IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "metal::dfdx"}, - IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32, - "metal::dfdx"}, - IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "metal::dfdx"}, - IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "metal::dfdy"}, - IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32, - "metal::dfdy"}, - IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "metal::dfdy"}, - IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "metal::exp"}, - IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "metal::exp2"}, + "determinant"}, + IntrinsicData{IntrinsicType::kDistance, ParamType::kF32, "distance"}, + IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "dot"}, + IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "dfdx"}, + IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32, "dfdx"}, + IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "dfdx"}, + IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "dfdy"}, + IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32, "dfdy"}, + IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "dfdy"}, + IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "exp"}, + IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "exp2"}, IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32, - "metal::faceforward"}, - IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "metal::floor"}, - IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "metal::fma"}, - IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "metal::fract"}, - IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "metal::fwidth"}, - IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32, - "metal::fwidth"}, - IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32, - "metal::fwidth"}, - IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32, - "metal::rsqrt"}, - IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32, - "metal::isfinite"}, - IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "metal::isinf"}, - IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "metal::isnan"}, - IntrinsicData{IntrinsicType::kIsNormal, ParamType::kF32, - "metal::isnormal"}, - IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "metal::ldexp"}, - IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "metal::length"}, - IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "metal::log"}, - IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "metal::log2"}, - IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "metal::fmax"}, - IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "metal::max"}, - IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "metal::fmin"}, - IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "metal::min"}, - IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32, - "metal::normalize"}, + "faceforward"}, + IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "floor"}, + IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "fma"}, + IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "fract"}, + IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "fwidth"}, + IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32, "fwidth"}, + IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32, "fwidth"}, + IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32, "rsqrt"}, + IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32, "isfinite"}, + IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "isinf"}, + IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "isnan"}, + IntrinsicData{IntrinsicType::kIsNormal, ParamType::kF32, "isnormal"}, + IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "ldexp"}, + IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "length"}, + IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "log"}, + IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "log2"}, + IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "fmax"}, + IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "max"}, + IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "fmin"}, + IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "min"}, + IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32, "normalize"}, IntrinsicData{IntrinsicType::kPack4x8Snorm, ParamType::kF32, - "metal::pack_float_to_snorm4x8"}, + "pack_float_to_snorm4x8"}, IntrinsicData{IntrinsicType::kPack4x8Unorm, ParamType::kF32, - "metal::pack_float_to_unorm4x8"}, + "pack_float_to_unorm4x8"}, IntrinsicData{IntrinsicType::kPack2x16Snorm, ParamType::kF32, - "metal::pack_float_to_snorm2x16"}, + "pack_float_to_snorm2x16"}, IntrinsicData{IntrinsicType::kPack2x16Unorm, ParamType::kF32, - "metal::pack_float_to_unorm2x16"}, - IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "metal::pow"}, - IntrinsicData{IntrinsicType::kReflect, ParamType::kF32, - "metal::reflect"}, + "pack_float_to_unorm2x16"}, + IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "pow"}, + IntrinsicData{IntrinsicType::kReflect, ParamType::kF32, "reflect"}, IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32, - "metal::reverse_bits"}, - IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "metal::rint"}, - IntrinsicData{IntrinsicType::kSelect, ParamType::kF32, "metal::select"}, - IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "metal::sign"}, - IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "metal::sin"}, - IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "metal::sinh"}, + "reverse_bits"}, + IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "rint"}, + IntrinsicData{IntrinsicType::kSelect, ParamType::kF32, "select"}, + IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "sign"}, + IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "sin"}, + IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "sinh"}, IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32, - "metal::smoothstep"}, - IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "metal::sqrt"}, - IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "metal::step"}, - IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "metal::tan"}, - IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "metal::tanh"}, - IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "metal::trunc"}, + "smoothstep"}, + IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "sqrt"}, + IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "step"}, + IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "tan"}, + IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "tanh"}, + IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "trunc"}, IntrinsicData{IntrinsicType::kUnpack4x8Snorm, ParamType::kU32, - "metal::unpack_snorm4x8_to_float"}, + "unpack_snorm4x8_to_float"}, IntrinsicData{IntrinsicType::kUnpack4x8Unorm, ParamType::kU32, - "metal::unpack_unorm4x8_to_float"}, + "unpack_unorm4x8_to_float"}, IntrinsicData{IntrinsicType::kUnpack2x16Snorm, ParamType::kU32, - "metal::unpack_snorm2x16_to_float"}, + "unpack_snorm2x16_to_float"}, IntrinsicData{IntrinsicType::kUnpack2x16Unorm, ParamType::kU32, - "metal::unpack_unorm2x16_to_float"})); + "unpack_unorm2x16_to_float"})); TEST_F(MslGeneratorImplTest, Intrinsic_Call) { Global("param1", ty.vec2<f32>(), ast::StorageClass::kFunction); @@ -303,7 +293,7 @@ TEST_F(MslGeneratorImplTest, Intrinsic_Call) { gen.increment_indent(); ASSERT_TRUE(gen.EmitExpression(call)) << gen.error(); - EXPECT_EQ(gen.result(), " metal::dot(param1, param2)"); + EXPECT_EQ(gen.result(), " dot(param1, param2)"); } TEST_F(MslGeneratorImplTest, Pack2x16Float) { diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc index 36067fc084..1a6d5a219f 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc @@ -58,6 +58,7 @@ TEST_F(MslGeneratorImplTest, Generate) { ASSERT_TRUE(gen.Generate()) << gen.error(); EXPECT_EQ(gen.result(), R"(#include <metal_stdlib> +using namespace metal; kernel void my_func() { return; }