writer/msl: Emit texture builtin functions
Bug: tint:145
Change-Id: I8a2d10c2e7239c81a11933c009d9175d4f4d8577
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/33782
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/BUILD.gn b/BUILD.gn
index bf1d393..cd9c72a 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -1114,6 +1114,7 @@
"src/writer/msl/generator_impl_if_test.cc",
"src/writer/msl/generator_impl_import_test.cc",
"src/writer/msl/generator_impl_intrinsic_test.cc",
+ "src/writer/msl/generator_impl_intrinsic_texture_test.cc",
"src/writer/msl/generator_impl_loop_test.cc",
"src/writer/msl/generator_impl_member_accessor_test.cc",
"src/writer/msl/generator_impl_module_constant_test.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index e41d3b7..45b2c88 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -657,6 +657,7 @@
writer/msl/generator_impl_if_test.cc
writer/msl/generator_impl_import_test.cc
writer/msl/generator_impl_intrinsic_test.cc
+ writer/msl/generator_impl_intrinsic_texture_test.cc
writer/msl/generator_impl_loop_test.cc
writer/msl/generator_impl_member_accessor_test.cc
writer/msl/generator_impl_module_constant_test.cc
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index f0a1b8d..a7fe323 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -545,8 +545,7 @@
auto name = generate_intrinsic_name(ident->intrinsic());
if (name.empty()) {
if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
- error_ = "Textures not implemented yet";
- return false;
+ return EmitTextureCall(expr);
}
name = generate_builtin_name(ident);
if (name.empty()) {
@@ -654,6 +653,107 @@
return true;
}
+bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
+ auto* ident = expr->func()->AsIdentifier();
+
+ auto params = expr->params();
+ auto* signature = static_cast<const ast::intrinsic::TextureSignature*>(
+ ident->intrinsic_signature());
+ auto& pidx = signature->params.idx;
+ auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
+
+ if (!EmitExpression(params[pidx.texture]))
+ return false;
+
+ switch (ident->intrinsic()) {
+ case ast::Intrinsic::kTextureSample:
+ case ast::Intrinsic::kTextureSampleBias:
+ case ast::Intrinsic::kTextureSampleLevel:
+ case ast::Intrinsic::kTextureSampleGrad:
+ out_ << ".sample(";
+ break;
+ case ast::Intrinsic::kTextureSampleCompare:
+ out_ << ".sample_compare(";
+ break;
+ default:
+ error_ = "Internal compiler error: Unhandled texture intrinsic '" +
+ ident->name() + "'";
+ break;
+ }
+
+ if (!EmitExpression(params[pidx.sampler])) {
+ return false;
+ }
+
+ for (auto idx : {pidx.coords, pidx.array_index, pidx.depth_ref}) {
+ if (idx != kNotUsed) {
+ out_ << ", ";
+ if (!EmitExpression(params[idx]))
+ return false;
+ }
+ }
+
+ if (pidx.bias != kNotUsed) {
+ out_ << ", bias(";
+ if (!EmitExpression(params[pidx.bias])) {
+ return false;
+ }
+ out_ << ")";
+ }
+ if (pidx.level != kNotUsed) {
+ out_ << ", level(";
+ if (!EmitExpression(params[pidx.level])) {
+ return false;
+ }
+ out_ << ")";
+ }
+ if (pidx.ddx != kNotUsed) {
+ auto dim = params[pidx.texture]
+ ->result_type()
+ ->UnwrapPtrIfNeeded()
+ ->AsTexture()
+ ->dim();
+ switch (dim) {
+ case ast::type::TextureDimension::k2d:
+ case ast::type::TextureDimension::k2dArray:
+ out_ << ", gradient2d(";
+ break;
+ case ast::type::TextureDimension::k3d:
+ out_ << ", gradient3d(";
+ break;
+ case ast::type::TextureDimension::kCube:
+ case ast::type::TextureDimension::kCubeArray:
+ out_ << ", gradientcube(";
+ break;
+ default: {
+ std::stringstream err;
+ err << "MSL does not support gradients for " << dim << " textures";
+ error_ = err.str();
+ return false;
+ }
+ }
+ if (!EmitExpression(params[pidx.ddx])) {
+ return false;
+ }
+ out_ << ", ";
+ if (!EmitExpression(params[pidx.ddy])) {
+ return false;
+ }
+ out_ << ")";
+ }
+
+ if (pidx.offset != kNotUsed) {
+ out_ << ", ";
+ if (!EmitExpression(params[pidx.offset])) {
+ return false;
+ }
+ }
+
+ out_ << ")";
+
+ return true;
+}
+
std::string GeneratorImpl::generate_builtin_name(
ast::IdentifierExpression* ident) {
std::string out = "metal::";
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 0c866af..7cbd5d6 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -95,6 +95,11 @@
/// @param expr the call expression
/// @returns true if the call expression is emitted
bool EmitCall(ast::CallExpression* expr);
+ /// Handles generating a call to a texture function (`textureSample`,
+ /// `textureSampleGrad`, etc)
+ /// @param expr the call expression
+ /// @returns true if the call expression is emitted
+ bool EmitTextureCall(ast::CallExpression* expr);
/// Handles a case statement
/// @param stmt the statement
/// @returns true if the statement was emitted successfully
diff --git a/src/writer/msl/generator_impl_intrinsic_texture_test.cc b/src/writer/msl/generator_impl_intrinsic_texture_test.cc
new file mode 100644
index 0000000..c8079da
--- /dev/null
+++ b/src/writer/msl/generator_impl_intrinsic_texture_test.cc
@@ -0,0 +1,207 @@
+// 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 "gtest/gtest.h"
+#include "src/ast/builder.h"
+#include "src/ast/intrinsic_texture_helper_test.h"
+#include "src/ast/type/depth_texture_type.h"
+#include "src/ast/type/sampled_texture_type.h"
+#include "src/context.h"
+#include "src/type_determiner.h"
+#include "src/writer/msl/generator_impl.h"
+
+namespace tint {
+namespace writer {
+namespace msl {
+namespace {
+
+std::string expected_texture_overload(
+ ast::intrinsic::test::ValidTextureOverload overload) {
+ using ValidTextureOverload = ast::intrinsic::test::ValidTextureOverload;
+ switch (overload) {
+ case ValidTextureOverload::kSample1dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, 1.0f))";
+ case ValidTextureOverload::kSample1dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, 1.0f, 2u))";
+ case ValidTextureOverload::kSample2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f)))";
+ case ValidTextureOverload::kSample2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), int2(3, 4)))";
+ case ValidTextureOverload::kSample2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u))";
+ case ValidTextureOverload::kSample2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, int2(4, 5)))";
+ case ValidTextureOverload::kSample3dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f)))";
+ case ValidTextureOverload::kSample3dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), int3(4, 5, 6)))";
+ case ValidTextureOverload::kSampleCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f)))";
+ case ValidTextureOverload::kSampleCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u))";
+ case ValidTextureOverload::kSampleDepth2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f)))";
+ case ValidTextureOverload::kSampleDepth2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), int2(3, 4)))";
+ case ValidTextureOverload::kSampleDepth2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u))";
+ case ValidTextureOverload::kSampleDepth2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, int2(4, 5)))";
+ case ValidTextureOverload::kSampleDepthCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f)))";
+ case ValidTextureOverload::kSampleDepthCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u))";
+ case ValidTextureOverload::kSampleBias2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), bias(3.0f)))";
+ case ValidTextureOverload::kSampleBias2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), bias(3.0f), int2(4, 5)))";
+ case ValidTextureOverload::kSampleBias2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 4u, bias(3.0f)))";
+ case ValidTextureOverload::kSampleBias2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, bias(4.0f), int2(5, 6)))";
+ case ValidTextureOverload::kSampleBias3dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), bias(4.0f)))";
+ case ValidTextureOverload::kSampleBias3dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), bias(4.0f), int3(5, 6, 7)))";
+ case ValidTextureOverload::kSampleBiasCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), bias(4.0f)))";
+ case ValidTextureOverload::kSampleBiasCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 3u, bias(4.0f)))";
+ case ValidTextureOverload::kSampleLevel2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), level(3.0f)))";
+ case ValidTextureOverload::kSampleLevel2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), level(3.0f), int2(4, 5)))";
+ case ValidTextureOverload::kSampleLevel2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, level(4.0f)))";
+ case ValidTextureOverload::kSampleLevel2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, level(4.0f), int2(5, 6)))";
+ case ValidTextureOverload::kSampleLevel3dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), level(4.0f)))";
+ case ValidTextureOverload::kSampleLevel3dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), level(4.0f), int3(5, 6, 7)))";
+ case ValidTextureOverload::kSampleLevelCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), level(4.0f)))";
+ case ValidTextureOverload::kSampleLevelCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u, level(5.0f)))";
+ case ValidTextureOverload::kSampleLevelDepth2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), level(3u)))";
+ case ValidTextureOverload::kSampleLevelDepth2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), level(3u), int2(4, 5)))";
+ case ValidTextureOverload::kSampleLevelDepth2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, level(4u)))";
+ case ValidTextureOverload::kSampleLevelDepth2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, level(4u), int2(5, 6)))";
+ case ValidTextureOverload::kSampleLevelDepthCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), level(4u)))";
+ case ValidTextureOverload::kSampleLevelDepthCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u, level(5u)))";
+ case ValidTextureOverload::kSampleGrad2dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), gradient2d(float2(3.0f, 4.0f), float2(5.0f, 6.0f))))";
+ case ValidTextureOverload::kSampleGrad2dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), gradient2d(float2(3.0f, 4.0f), float2(5.0f, 6.0f)), int2(7, 8)))";
+ case ValidTextureOverload::kSampleGrad2dArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, gradient2d(float2(4.0f, 5.0f), float2(6.0f, 7.0f))))";
+ case ValidTextureOverload::kSampleGrad2dArrayOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float2(1.0f, 2.0f), 3u, gradient2d(float2(4.0f, 5.0f), float2(6.0f, 7.0f)), int2(8, 9)))";
+ case ValidTextureOverload::kSampleGrad3dF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), gradient3d(float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f))))";
+ case ValidTextureOverload::kSampleGrad3dOffsetF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), gradient3d(float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f)), int3(10, 11, 12)))";
+ case ValidTextureOverload::kSampleGradCubeF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), gradientcube(float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f))))";
+ case ValidTextureOverload::kSampleGradCubeArrayF32:
+ return R"(texture_tint_0.sample(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u, gradientcube(float3(5.0f, 6.0f, 7.0f), float3(8.0f, 9.0f, 10.0f))))";
+ case ValidTextureOverload::kSampleGradDepth2dF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float2(1.0f, 2.0f), 3.0f))";
+ case ValidTextureOverload::kSampleGradDepth2dOffsetF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float2(1.0f, 2.0f), 3.0f, int2(4, 5)))";
+ case ValidTextureOverload::kSampleGradDepth2dArrayF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float2(1.0f, 2.0f), 4u, 3.0f))";
+ case ValidTextureOverload::kSampleGradDepth2dArrayOffsetF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float2(1.0f, 2.0f), 4u, 3.0f, int2(5, 6)))";
+ case ValidTextureOverload::kSampleGradDepthCubeF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4.0f))";
+ case ValidTextureOverload::kSampleGradDepthCubeArrayF32:
+ return R"(texture_tint_0.sample_compare(sampler_tint_0, float3(1.0f, 2.0f, 3.0f), 4u, 5.0f))";
+ }
+ return "<unmatched texture overload>";
+} // LINT - Ignore the length of this function
+
+class MslGeneratorIntrinsicTextureTest
+ : public ast::BuilderWithContextAndModule,
+ public testing::TestWithParam<ast::intrinsic::test::TextureOverloadCase> {
+ protected:
+ void OnVariableBuilt(ast::Variable* var) override {
+ td.RegisterVariableForTesting(var);
+ }
+
+ /// The type determiner
+ TypeDeterminer td{ctx, mod};
+ /// The generator
+ GeneratorImpl gen{ctx, mod};
+};
+
+TEST_P(MslGeneratorIntrinsicTextureTest, Call) {
+ auto param = GetParam();
+
+ ast::type::Type* datatype = nullptr;
+ switch (param.texture_data_type) {
+ case ast::intrinsic::test::TextureDataType::kF32:
+ datatype = ty.f32;
+ break;
+ case ast::intrinsic::test::TextureDataType::kU32:
+ datatype = ty.u32;
+ break;
+ case ast::intrinsic::test::TextureDataType::kI32:
+ datatype = ty.i32;
+ break;
+ }
+
+ ast::type::SamplerType sampler_type{param.sampler_kind};
+ switch (param.texture_kind) {
+ case ast::intrinsic::test::TextureKind::kRegular:
+ Var("texture", ast::StorageClass::kNone,
+ mod->create<ast::type::SampledTextureType>(param.texture_dimension,
+ datatype));
+ break;
+
+ case ast::intrinsic::test::TextureKind::kDepth:
+ Var("texture", ast::StorageClass::kNone,
+ mod->create<ast::type::DepthTextureType>(param.texture_dimension));
+ break;
+ }
+
+ Var("sampler", ast::StorageClass::kNone, &sampler_type);
+
+ ast::CallExpression call{Expr(param.function), param.args(this)};
+
+ EXPECT_TRUE(td.DetermineResultType(&call)) << td.error();
+
+ ASSERT_TRUE(gen.EmitExpression(&call)) << gen.error();
+
+ auto expected = expected_texture_overload(param.overload);
+ EXPECT_EQ(gen.result(), expected);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ MslGeneratorIntrinsicTextureTest,
+ MslGeneratorIntrinsicTextureTest,
+ testing::ValuesIn(ast::intrinsic::test::TextureOverloadCase::ValidCases()));
+
+} // namespace
+} // namespace msl
+} // namespace writer
+} // namespace tint
