[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>
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index a603d07..dee90cb 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -122,6 +122,7 @@
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 @@
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 69f2ced..714c0b2 100644
--- a/src/writer/msl/generator_impl_function_test.cc
+++ b/src/writer/msl/generator_impl_function_test.cc
@@ -69,6 +69,7 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 4894db3..5176a92 100644
--- a/src/writer/msl/generator_impl_import_test.cc
+++ b/src/writer/msl/generator_impl_import_test.cc
@@ -65,8 +65,7 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 cd28dce..7f26363 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -199,98 +199,88 @@
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 @@
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 36067fc..1a6d5a2 100644
--- a/src/writer/msl/generator_impl_test.cc
+++ b/src/writer/msl/generator_impl_test.cc
@@ -58,6 +58,7 @@
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
+using namespace metal;
kernel void my_func() {
return;
}
