[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

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

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;
 };

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

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) {

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;
 }