Add option to auto generate bindings for external textures
With this change, the backend sanitizers always run the
MultiplanarExternalTexture transform. If the new option is enabled, it
auto-generates bindings for this transform.
This change also enables this auto-generation for the Tint commandline
application, as well as for the fuzzers.
Bug: chromium:1310623
Change-Id: I3c661c4753dc67c0212051d09024cbeda3939f8c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/85542
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 6c525c8..58ec6a4 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -527,6 +527,8 @@
"writer/array_length_from_uniform_options.h",
"writer/float_to_string.cc",
"writer/float_to_string.h",
+ "writer/generate_external_texture_bindings.cc",
+ "writer/generate_external_texture_bindings.h",
"writer/text.cc",
"writer/text.h",
"writer/text_generator.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 207b988..138dfde 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -451,6 +451,8 @@
writer/array_length_from_uniform_options.h
writer/float_to_string.cc
writer/float_to_string.h
+ writer/generate_external_texture_bindings.cc
+ writer/generate_external_texture_bindings.h
writer/text_generator.cc
writer/text_generator.h
writer/text.cc
@@ -810,6 +812,7 @@
utils/unique_vector_test.cc
writer/append_vector_test.cc
writer/float_to_string_test.cc
+ writer/generate_external_texture_bindings_test.cc
writer/text_generator_test.cc
)
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index ff3587c..8a0d293 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -569,6 +569,7 @@
// TODO(jrprice): Provide a way for the user to set non-default options.
tint::writer::spirv::Options gen_options;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
+ gen_options.generate_external_texture_bindings = true;
auto result = tint::writer::spirv::Generate(program, gen_options);
if (!result.success) {
PrintWGSL(std::cerr, *program);
@@ -717,6 +718,7 @@
// TODO(jrprice): Provide a way for the user to set non-default options.
tint::writer::msl::Options gen_options;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
+ gen_options.generate_external_texture_bindings = true;
auto result = tint::writer::msl::Generate(input_program, gen_options);
if (!result.success) {
PrintWGSL(std::cerr, *program);
@@ -771,6 +773,7 @@
// TODO(jrprice): Provide a way for the user to set non-default options.
tint::writer::hlsl::Options gen_options;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
+ gen_options.generate_external_texture_bindings = true;
auto result = tint::writer::hlsl::Generate(program, gen_options);
if (!result.success) {
PrintWGSL(std::cerr, *program);
@@ -846,6 +849,7 @@
auto generate = [&](const tint::Program* prg,
const std::string entry_point_name) -> bool {
tint::writer::glsl::Options gen_options;
+ gen_options.generate_external_texture_bindings = true;
auto result =
tint::writer::glsl::Generate(prg, gen_options, entry_point_name);
if (!result.success) {
diff --git a/src/tint/fuzzers/data_builder.h b/src/tint/fuzzers/data_builder.h
index 31a1a98..e0c104a 100644
--- a/src/tint/fuzzers/data_builder.h
+++ b/src/tint/fuzzers/data_builder.h
@@ -172,6 +172,7 @@
b->build(out.fixed_sample_mask);
b->build(out.emit_vertex_point_size);
b->build(out.disable_workgroup_init);
+ b->build(out.generate_external_texture_bindings);
b->build(out.array_length_from_uniform);
return out;
}
diff --git a/src/tint/transform/glsl.cc b/src/tint/transform/glsl.cc
index 5bb0006..df6b6b3 100644
--- a/src/tint/transform/glsl.cc
+++ b/src/tint/transform/glsl.cc
@@ -37,6 +37,7 @@
#include "src/tint/transform/unshadow.h"
#include "src/tint/transform/unwind_discard_functions.h"
#include "src/tint/transform/zero_init_workgroup_memory.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl);
TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl::Config);
@@ -91,6 +92,13 @@
manager.Add<SimplifyPointers>();
manager.Add<RemovePhonies>();
+
+ if (cfg && cfg->generate_external_texture_bindings) {
+ auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
+ data.Add<MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+ }
+ manager.Add<MultiplanarExternalTexture>();
+
manager.Add<CombineSamplers>();
if (auto* binding_info = inputs.Get<CombineSamplers::BindingInfo>()) {
data.Add<CombineSamplers::BindingInfo>(*binding_info);
@@ -106,6 +114,7 @@
BindingRemapper::AccessControls ac;
data.Add<BindingRemapper::Remappings>(bp, ac, /* mayCollide */ true);
}
+
manager.Add<PromoteInitializersToConstVar>();
manager.Add<AddEmptyEntryPoint>();
@@ -124,8 +133,13 @@
return Output{Program(std::move(builder))};
}
-Glsl::Config::Config(const std::string& ep, bool disable_wi)
- : entry_point(ep), disable_workgroup_init(disable_wi) {}
+Glsl::Config::Config(const std::string& entry_point_in,
+ bool disable_workgroup_init_in,
+ bool generate_external_texture_bindings_in)
+ : entry_point(entry_point_in),
+ disable_workgroup_init(disable_workgroup_init_in),
+ generate_external_texture_bindings(
+ generate_external_texture_bindings_in) {}
Glsl::Config::Config(const Config&) = default;
Glsl::Config::~Config() = default;
diff --git a/src/tint/transform/glsl.h b/src/tint/transform/glsl.h
index baec29a..414f871 100644
--- a/src/tint/transform/glsl.h
+++ b/src/tint/transform/glsl.h
@@ -37,8 +37,11 @@
/// @param entry_point the root entry point function to generate
/// @param disable_workgroup_init `true` to disable workgroup memory zero
/// initialization
+ /// @param generate_external_texture_bindings 'true' to generates binding
+ /// mappings for external textures
explicit Config(const std::string& entry_point,
- bool disable_workgroup_init = false);
+ bool disable_workgroup_init,
+ bool generate_external_texture_bindings);
/// Copy constructor
Config(const Config&);
@@ -51,6 +54,9 @@
/// Set to `true` to disable workgroup memory zero initialization
bool disable_workgroup_init = false;
+
+ /// Set to 'true' to generates binding mappings for external textures
+ bool generate_external_texture_bindings = false;
};
/// Constructor
diff --git a/src/tint/writer/generate_external_texture_bindings.cc b/src/tint/writer/generate_external_texture_bindings.cc
new file mode 100644
index 0000000..44bd262
--- /dev/null
+++ b/src/tint/writer/generate_external_texture_bindings.cc
@@ -0,0 +1,59 @@
+// Copyright 2022 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 "src/tint/writer/generate_external_texture_bindings.h"
+
+#include <algorithm>
+#include <unordered_map>
+#include <vector>
+
+#include "src/tint/ast/external_texture.h"
+#include "src/tint/ast/module.h"
+#include "src/tint/program.h"
+#include "src/tint/sem/external_texture_type.h"
+#include "src/tint/sem/variable.h"
+
+namespace tint::writer {
+
+transform::MultiplanarExternalTexture::BindingsMap
+GenerateExternalTextureBindings(const Program* program) {
+ // TODO(tint:1491): Use Inspector once we can get binding info for all
+ // variables, not just those referenced by entry points.
+
+ // Collect next valid binding number per group
+ std::unordered_map<uint32_t, uint32_t> group_to_next_binding_number;
+ std::vector<sem::BindingPoint> ext_tex_bps;
+ for (auto* var : program->AST().GlobalVariables()) {
+ if (auto* sem_var = program->Sem().Get(var)->As<sem::GlobalVariable>()) {
+ auto bp = sem_var->BindingPoint();
+ auto& n = group_to_next_binding_number[bp.group];
+ n = std::max(n, bp.binding + 1);
+
+ if (sem_var->Type()->UnwrapRef()->Is<sem::ExternalTexture>()) {
+ ext_tex_bps.emplace_back(bp);
+ }
+ }
+ }
+
+ transform::MultiplanarExternalTexture::BindingsMap new_bindings_map;
+ for (auto bp : ext_tex_bps) {
+ uint32_t g = bp.group;
+ uint32_t& next_num = group_to_next_binding_number[g];
+ auto new_bps = transform::BindingPoints{{g, next_num++}, {g, next_num++}};
+ new_bindings_map[bp] = new_bps;
+ }
+ return new_bindings_map;
+}
+
+} // namespace tint::writer
diff --git a/src/tint/writer/generate_external_texture_bindings.h b/src/tint/writer/generate_external_texture_bindings.h
new file mode 100644
index 0000000..6f321d5
--- /dev/null
+++ b/src/tint/writer/generate_external_texture_bindings.h
@@ -0,0 +1,27 @@
+// Copyright 2022 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.
+
+#ifndef SRC_TINT_WRITER_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
+#define SRC_TINT_WRITER_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
+
+#include "src/tint/transform/multiplanar_external_texture.h"
+
+namespace tint::writer {
+
+transform::MultiplanarExternalTexture::BindingsMap
+GenerateExternalTextureBindings(const Program* program);
+
+} // namespace tint::writer
+
+#endif // SRC_TINT_WRITER_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
diff --git a/src/tint/writer/generate_external_texture_bindings_test.cc b/src/tint/writer/generate_external_texture_bindings_test.cc
new file mode 100644
index 0000000..3163007
--- /dev/null
+++ b/src/tint/writer/generate_external_texture_bindings_test.cc
@@ -0,0 +1,131 @@
+// Copyright 2022 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 <utility>
+
+#include "gtest/gtest.h"
+#include "src/tint/program_builder.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
+
+namespace tint::writer {
+namespace {
+
+constexpr auto kUniform = ast::StorageClass::kUniform;
+
+class GenerateExternalTextureBindingsTest : public ::testing::Test {};
+
+TEST_F(GenerateExternalTextureBindingsTest, None) {
+ ProgramBuilder b;
+ b.WrapInFunction();
+
+ tint::Program program(std::move(b));
+ ASSERT_TRUE(program.IsValid());
+ auto bindings = GenerateExternalTextureBindings(&program);
+ ASSERT_TRUE(bindings.empty());
+}
+
+TEST_F(GenerateExternalTextureBindingsTest, One) {
+ ProgramBuilder b;
+ b.Global("v0", b.ty.external_texture(), b.GroupAndBinding(0, 0));
+ b.WrapInFunction();
+
+ tint::Program program(std::move(b));
+ ASSERT_TRUE(program.IsValid());
+ auto bindings = GenerateExternalTextureBindings(&program);
+ ASSERT_EQ(bindings.size(), 1u);
+
+ auto to = bindings[transform::BindingPoint{0, 0}];
+ ASSERT_EQ(to.plane_1.group, 0u);
+ ASSERT_EQ(to.params.group, 0u);
+ ASSERT_EQ(to.plane_1.binding, 1u);
+ ASSERT_EQ(to.params.binding, 2u);
+}
+
+TEST_F(GenerateExternalTextureBindingsTest, Two_SameGroup) {
+ ProgramBuilder b;
+ b.Global("v0", b.ty.external_texture(), b.GroupAndBinding(0, 0));
+ b.Global("v1", b.ty.external_texture(), b.GroupAndBinding(0, 1));
+ b.WrapInFunction();
+
+ tint::Program program(std::move(b));
+ ASSERT_TRUE(program.IsValid());
+ auto bindings = GenerateExternalTextureBindings(&program);
+ ASSERT_EQ(bindings.size(), 2u);
+
+ auto to0 = bindings[transform::BindingPoint{0, 0}];
+ ASSERT_EQ(to0.plane_1.group, 0u);
+ ASSERT_EQ(to0.params.group, 0u);
+ ASSERT_EQ(to0.plane_1.binding, 2u);
+ ASSERT_EQ(to0.params.binding, 3u);
+
+ auto to1 = bindings[transform::BindingPoint{0, 1}];
+ ASSERT_EQ(to1.plane_1.group, 0u);
+ ASSERT_EQ(to1.params.group, 0u);
+ ASSERT_EQ(to1.plane_1.binding, 4u);
+ ASSERT_EQ(to1.params.binding, 5u);
+}
+
+TEST_F(GenerateExternalTextureBindingsTest, Two_DifferentGroup) {
+ ProgramBuilder b;
+ b.Global("v0", b.ty.external_texture(), b.GroupAndBinding(0, 0));
+ b.Global("v1", b.ty.external_texture(), b.GroupAndBinding(1, 0));
+ b.WrapInFunction();
+
+ tint::Program program(std::move(b));
+ ASSERT_TRUE(program.IsValid());
+ auto bindings = GenerateExternalTextureBindings(&program);
+ ASSERT_EQ(bindings.size(), 2u);
+
+ auto to0 = bindings[transform::BindingPoint{0, 0}];
+ ASSERT_EQ(to0.plane_1.group, 0u);
+ ASSERT_EQ(to0.params.group, 0u);
+ ASSERT_EQ(to0.plane_1.binding, 1u);
+ ASSERT_EQ(to0.params.binding, 2u);
+
+ auto to1 = bindings[transform::BindingPoint{1, 0}];
+ ASSERT_EQ(to1.plane_1.group, 1u);
+ ASSERT_EQ(to1.params.group, 1u);
+ ASSERT_EQ(to1.plane_1.binding, 1u);
+ ASSERT_EQ(to1.params.binding, 2u);
+}
+
+TEST_F(GenerateExternalTextureBindingsTest, Two_WithOtherBindingsInSameGroup) {
+ ProgramBuilder b;
+ b.Global("v0", b.ty.i32(), b.GroupAndBinding(0, 0), kUniform);
+ b.Global("v1", b.ty.external_texture(), b.GroupAndBinding(0, 1));
+ b.Global("v2", b.ty.i32(), b.GroupAndBinding(0, 2), kUniform);
+ b.Global("v3", b.ty.external_texture(), b.GroupAndBinding(0, 3));
+ b.Global("v4", b.ty.i32(), b.GroupAndBinding(0, 4), kUniform);
+ b.WrapInFunction();
+
+ tint::Program program(std::move(b));
+ ASSERT_TRUE(program.IsValid()) << program.Diagnostics().str();
+ auto bindings = GenerateExternalTextureBindings(&program);
+ ASSERT_EQ(bindings.size(), 2u);
+
+ auto to0 = bindings[transform::BindingPoint{0, 1}];
+ ASSERT_EQ(to0.plane_1.group, 0u);
+ ASSERT_EQ(to0.params.group, 0u);
+ ASSERT_EQ(to0.plane_1.binding, 5u);
+ ASSERT_EQ(to0.params.binding, 6u);
+
+ auto to1 = bindings[transform::BindingPoint{0, 3}];
+ ASSERT_EQ(to1.plane_1.group, 0u);
+ ASSERT_EQ(to1.params.group, 0u);
+ ASSERT_EQ(to1.plane_1.binding, 7u);
+ ASSERT_EQ(to1.params.binding, 8u);
+}
+
+} // namespace
+} // namespace tint::writer
diff --git a/src/tint/writer/glsl/generator.cc b/src/tint/writer/glsl/generator.cc
index ab7079f..9ae88ad 100644
--- a/src/tint/writer/glsl/generator.cc
+++ b/src/tint/writer/glsl/generator.cc
@@ -43,7 +43,9 @@
options.allow_collisions);
data.Add<transform::CombineSamplers::BindingInfo>(
options.binding_map, options.placeholder_binding_point);
- data.Add<transform::Glsl::Config>(entry_point);
+ data.Add<transform::Glsl::Config>(entry_point,
+ /* disable_workgroup_init */ false,
+ options.generate_external_texture_bindings);
transform::Glsl sanitizer;
auto output = sanitizer.Run(program, data);
if (!output.program.IsValid()) {
diff --git a/src/tint/writer/glsl/generator.h b/src/tint/writer/glsl/generator.h
index 08b3e93..88be429 100644
--- a/src/tint/writer/glsl/generator.h
+++ b/src/tint/writer/glsl/generator.h
@@ -71,6 +71,9 @@
/// generated by the BindingRemapper transform
bool allow_collisions = false;
+ /// Set to 'true' to generates binding mappings for external textures
+ bool generate_external_texture_bindings = false;
+
/// The GLSL version to emit
Version version;
};
diff --git a/src/tint/writer/hlsl/generator.cc b/src/tint/writer/hlsl/generator.cc
index 536591f..b6aa19c 100644
--- a/src/tint/writer/hlsl/generator.cc
+++ b/src/tint/writer/hlsl/generator.cc
@@ -35,6 +35,7 @@
// Sanitize the program.
auto sanitized_result = Sanitize(program, options.root_constant_binding_point,
options.disable_workgroup_init,
+ options.generate_external_texture_bindings,
options.array_length_from_uniform);
if (!sanitized_result.program.IsValid()) {
result.success = false;
diff --git a/src/tint/writer/hlsl/generator.h b/src/tint/writer/hlsl/generator.h
index 63df62e..aa7cfa0 100644
--- a/src/tint/writer/hlsl/generator.h
+++ b/src/tint/writer/hlsl/generator.h
@@ -53,6 +53,8 @@
sem::BindingPoint root_constant_binding_point;
/// Set to `true` to disable workgroup memory zero initialization
bool disable_workgroup_init = false;
+ /// Set to 'true' to generates binding mappings for external textures
+ bool generate_external_texture_bindings = false;
/// Options used to specify a mapping of binding points to indices into a UBO
/// from which to load buffer sizes.
ArrayLengthFromUniformOptions array_length_from_uniform = {};
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index a8981f9..fccbdcd 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -71,6 +71,7 @@
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/writer/append_vector.h"
#include "src/tint/writer/float_to_string.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint {
namespace writer {
@@ -139,6 +140,7 @@
const Program* in,
sem::BindingPoint root_constant_binding_point,
bool disable_workgroup_init,
+ bool generate_external_texture_bindings,
const ArrayLengthFromUniformOptions& array_length_from_uniform) {
transform::Manager manager;
transform::DataMap data;
@@ -163,6 +165,13 @@
array_length_from_uniform_cfg.bindpoint_to_size_index =
array_length_from_uniform.bindpoint_to_size_index;
+ if (generate_external_texture_bindings) {
+ auto new_bindings_map = GenerateExternalTextureBindings(in);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ new_bindings_map);
+ }
+ manager.Add<transform::MultiplanarExternalTexture>();
+
manager.Add<transform::Unshadow>();
// LocalizeStructArrayAssignment must come after:
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index e3bb1b3..84b533d 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -77,6 +77,7 @@
const Program* program,
sem::BindingPoint root_constant_binding_point = {},
bool disable_workgroup_init = false,
+ bool generate_external_texture_bindings = false,
const ArrayLengthFromUniformOptions& array_length_from_uniform = {});
/// Implementation class for HLSL generator
diff --git a/src/tint/writer/hlsl/test_helper.h b/src/tint/writer/hlsl/test_helper.h
index 5ef28a2..cac4645 100644
--- a/src/tint/writer/hlsl/test_helper.h
+++ b/src/tint/writer/hlsl/test_helper.h
@@ -78,9 +78,11 @@
<< formatter.format(program->Diagnostics());
}();
- auto sanitized_result = Sanitize(
- program.get(), options.root_constant_binding_point,
- options.disable_workgroup_init, options.array_length_from_uniform);
+ auto sanitized_result =
+ Sanitize(program.get(), options.root_constant_binding_point,
+ options.disable_workgroup_init,
+ options.generate_external_texture_bindings,
+ options.array_length_from_uniform);
[&]() {
ASSERT_TRUE(sanitized_result.program.IsValid())
<< formatter.format(sanitized_result.program.Diagnostics());
diff --git a/src/tint/writer/msl/generator.cc b/src/tint/writer/msl/generator.cc
index fa3c2f0..715af8d 100644
--- a/src/tint/writer/msl/generator.cc
+++ b/src/tint/writer/msl/generator.cc
@@ -38,6 +38,7 @@
auto sanitized_result = Sanitize(
program, options.buffer_size_ubo_index, options.fixed_sample_mask,
options.emit_vertex_point_size, options.disable_workgroup_init,
+ options.generate_external_texture_bindings,
options.array_length_from_uniform);
if (!sanitized_result.program.IsValid()) {
result.success = false;
diff --git a/src/tint/writer/msl/generator.h b/src/tint/writer/msl/generator.h
index aa7f8c5..27f9968 100644
--- a/src/tint/writer/msl/generator.h
+++ b/src/tint/writer/msl/generator.h
@@ -61,6 +61,9 @@
/// Set to `true` to disable workgroup memory zero initialization
bool disable_workgroup_init = false;
+ /// Set to 'true' to generates binding mappings for external textures
+ bool generate_external_texture_bindings = false;
+
/// Options used to specify a mapping of binding points to indices into a UBO
/// from which to load buffer sizes.
ArrayLengthFromUniformOptions array_length_from_uniform = {};
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 76d19ee..5807b70 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -77,6 +77,7 @@
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/writer/float_to_string.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint {
namespace writer {
@@ -113,6 +114,7 @@
private:
std::ostream& s;
};
+
} // namespace
SanitizedResult::SanitizedResult() = default;
@@ -125,6 +127,7 @@
uint32_t fixed_sample_mask,
bool emit_vertex_point_size,
bool disable_workgroup_init,
+ bool generate_external_texture_bindings,
const ArrayLengthFromUniformOptions& array_length_from_uniform) {
transform::Manager manager;
transform::DataMap data;
@@ -167,6 +170,13 @@
transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, fixed_sample_mask,
emit_vertex_point_size);
+ if (generate_external_texture_bindings) {
+ auto new_bindings_map = GenerateExternalTextureBindings(in);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ new_bindings_map);
+ }
+ manager.Add<transform::MultiplanarExternalTexture>();
+
manager.Add<transform::Unshadow>();
if (!disable_workgroup_init) {
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index 17ad9bd..c973fe6 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -84,6 +84,7 @@
uint32_t fixed_sample_mask = 0xFFFFFFFF,
bool emit_vertex_point_size = false,
bool disable_workgroup_init = false,
+ bool generate_external_texture_bindings = false,
const ArrayLengthFromUniformOptions& array_length_from_uniform = {});
/// Implementation class for MSL generator
diff --git a/src/tint/writer/msl/test_helper.h b/src/tint/writer/msl/test_helper.h
index e1d383d..2e0f3f1 100644
--- a/src/tint/writer/msl/test_helper.h
+++ b/src/tint/writer/msl/test_helper.h
@@ -79,6 +79,7 @@
auto result = Sanitize(
program.get(), options.buffer_size_ubo_index, options.fixed_sample_mask,
options.emit_vertex_point_size, options.disable_workgroup_init,
+ options.generate_external_texture_bindings,
options.array_length_from_uniform);
[&]() {
ASSERT_TRUE(result.program.IsValid())
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 52ba914..5de6e46 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -60,6 +60,7 @@
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/writer/append_vector.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint {
namespace writer {
@@ -258,7 +259,8 @@
SanitizedResult Sanitize(const Program* in,
bool emit_vertex_point_size,
- bool disable_workgroup_init) {
+ bool disable_workgroup_init,
+ bool generate_external_texture_bindings) {
transform::Manager manager;
transform::DataMap data;
@@ -275,6 +277,13 @@
manager.Add<transform::BuiltinPolyfill>();
}
+ if (generate_external_texture_bindings) {
+ auto new_bindings_map = GenerateExternalTextureBindings(in);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ new_bindings_map);
+ }
+ manager.Add<transform::MultiplanarExternalTexture>();
+
manager.Add<transform::Unshadow>();
if (!disable_workgroup_init) {
manager.Add<transform::ZeroInitWorkgroupMemory>();
diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h
index ef48afd..2a94635 100644
--- a/src/tint/writer/spirv/builder.h
+++ b/src/tint/writer/spirv/builder.h
@@ -65,7 +65,8 @@
/// @returns the sanitized program and any supplementary information
SanitizedResult Sanitize(const Program* program,
bool emit_vertex_point_size = false,
- bool disable_workgroup_init = false);
+ bool disable_workgroup_init = false,
+ bool generate_external_texture_bindings = false);
/// Builder class to create SPIR-V instructions from a module.
class Builder {
diff --git a/src/tint/writer/spirv/generator.cc b/src/tint/writer/spirv/generator.cc
index 4b3b72e..f27514a 100644
--- a/src/tint/writer/spirv/generator.cc
+++ b/src/tint/writer/spirv/generator.cc
@@ -32,7 +32,8 @@
options.disable_workgroup_init ||
options.use_zero_initialize_workgroup_memory_extension;
auto sanitized_result = Sanitize(program, options.emit_vertex_point_size,
- disable_workgroup_init_in_sanitizer);
+ disable_workgroup_init_in_sanitizer,
+ options.generate_external_texture_bindings);
if (!sanitized_result.program.IsValid()) {
result.success = false;
result.error = sanitized_result.program.Diagnostics().str();
diff --git a/src/tint/writer/spirv/generator.h b/src/tint/writer/spirv/generator.h
index 7cf9654..191e43a 100644
--- a/src/tint/writer/spirv/generator.h
+++ b/src/tint/writer/spirv/generator.h
@@ -42,6 +42,9 @@
/// Set to `true` to disable workgroup memory zero initialization
bool disable_workgroup_init = false;
+ /// Set to 'true' to generates binding mappings for external textures
+ bool generate_external_texture_bindings = false;
+
/// Set to `true` to initialize workgroup memory with OpConstantNull when
/// VK_KHR_zero_initialize_workgroup_memory is enabled.
bool use_zero_initialize_workgroup_memory_extension = false;
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index b6bcdd4..d2ca288 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -376,6 +376,7 @@
sources = [
"../../src/tint/writer/append_vector_test.cc",
"../../src/tint/writer/float_to_string_test.cc",
+ "../../src/tint/writer/generate_external_texture_bindings_test.cc",
"../../src/tint/writer/text_generator_test.cc",
]
}
diff --git a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.glsl b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.glsl
index 3808a80..b351a33 100644
--- a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.glsl
@@ -1,11 +1,99 @@
-SKIP: FAILED
+#version 310 es
-../../src/tint/writer/glsl/generator_impl.cc:2544 internal compiler error: Multiplanar external texture transform was not run.
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+uniform highp sampler2D arg_0_1;
+void textureDimensions_ba1481() {
+ ivec2 res = textureSize(arg_0_1, 0);
+}
+
+vec4 vertex_main() {
+ textureDimensions_ba1481();
+ return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+#version 310 es
+precision mediump float;
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+uniform highp sampler2D arg_0_1;
+void textureDimensions_ba1481() {
+ ivec2 res = textureSize(arg_0_1, 0);
+}
+
+void fragment_main() {
+ textureDimensions_ba1481();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+uniform highp sampler2D arg_0_1;
+void textureDimensions_ba1481() {
+ ivec2 res = textureSize(arg_0_1, 0);
+}
+
+void compute_main() {
+ textureDimensions_ba1481();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
diff --git a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.hlsl b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.hlsl
index b3cea20..4697e5b 100644
--- a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.hlsl
@@ -1,11 +1,38 @@
-SKIP: FAILED
+Texture2D<float4> ext_tex_plane_1 : register(t1, space1);
+cbuffer cbuffer_ext_tex_params : register(b2, space1) {
+ uint4 ext_tex_params[2];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
-C:\src\tint2\src\writer\hlsl\generator_impl.cc:3632 internal compiler error: Multiplanar external texture transform was not run.
+void textureDimensions_ba1481() {
+ int2 tint_tmp;
+ arg_0.GetDimensions(tint_tmp.x, tint_tmp.y);
+ int2 res = tint_tmp;
+}
+struct tint_symbol {
+ float4 value : SV_Position;
+};
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+float4 vertex_main_inner() {
+ textureDimensions_ba1481();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ textureDimensions_ba1481();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ textureDimensions_ba1481();
+ return;
+}
diff --git a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.msl b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.msl
index 21d485b..1faa96a 100644
--- a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.msl
+++ b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.msl
@@ -1,11 +1,41 @@
-SKIP: FAILED
+#include <metal_stdlib>
-C:\src\tint2\src\writer\msl\generator_impl.cc:2344 internal compiler error: Multiplanar external texture transform was not run.
+using namespace metal;
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+void textureDimensions_ba1481(texture2d<float, access::sample> tint_symbol_1) {
+ int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
+}
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_2) {
+ textureDimensions_ba1481(tint_symbol_2);
+ return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_3 [[texture(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_3);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_4 [[texture(0)]]) {
+ textureDimensions_ba1481(tint_symbol_4);
+ return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_5 [[texture(0)]]) {
+ textureDimensions_ba1481(tint_symbol_5);
+ return;
+}
+
diff --git a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm
index 87554bb..461515f 100644
--- a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm
@@ -1,11 +1,100 @@
-SKIP: FAILED
-
-C:\src\tint2\src\writer\spirv\builder.cc:4013 internal compiler error: Multiplanar external texture transform was not run.
-
-
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 43
+; Schema: 0
+ OpCapability Shader
+ OpCapability ImageQuery
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %ext_tex_plane_1 "ext_tex_plane_1"
+ OpName %ExternalTextureParams "ExternalTextureParams"
+ OpMemberName %ExternalTextureParams 0 "numPlanes"
+ OpMemberName %ExternalTextureParams 1 "vr"
+ OpMemberName %ExternalTextureParams 2 "ug"
+ OpMemberName %ExternalTextureParams 3 "vg"
+ OpMemberName %ExternalTextureParams 4 "ub"
+ OpName %ext_tex_params "ext_tex_params"
+ OpName %arg_0 "arg_0"
+ OpName %textureDimensions_ba1481 "textureDimensions_ba1481"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %ext_tex_plane_1 DescriptorSet 1
+ OpDecorate %ext_tex_plane_1 Binding 1
+ OpDecorate %ExternalTextureParams Block
+ OpMemberDecorate %ExternalTextureParams 0 Offset 0
+ OpMemberDecorate %ExternalTextureParams 1 Offset 4
+ OpMemberDecorate %ExternalTextureParams 2 Offset 8
+ OpMemberDecorate %ExternalTextureParams 3 Offset 12
+ OpMemberDecorate %ExternalTextureParams 4 Offset 16
+ OpDecorate %ext_tex_params NonWritable
+ OpDecorate %ext_tex_params DescriptorSet 1
+ OpDecorate %ext_tex_params Binding 2
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %uint = OpTypeInt 32 0
+%ExternalTextureParams = OpTypeStruct %uint %float %float %float %float
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%ext_tex_params = OpVariable %_ptr_Uniform_ExternalTextureParams Uniform
+ %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %void = OpTypeVoid
+ %17 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %int_0 = OpConstant %int 0
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+ %28 = OpConstantNull %v2int
+ %29 = OpTypeFunction %v4float
+ %float_1 = OpConstant %float 1
+%textureDimensions_ba1481 = OpFunction %void None %17
+ %20 = OpLabel
+ %res = OpVariable %_ptr_Function_v2int Function %28
+ %24 = OpLoad %11 %arg_0
+ %21 = OpImageQuerySizeLod %v2int %24 %int_0
+ OpStore %res %21
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %29
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %textureDimensions_ba1481
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %17
+ %34 = OpLabel
+ %35 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %35
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %17
+ %38 = OpLabel
+ %39 = OpFunctionCall %void %textureDimensions_ba1481
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %17
+ %41 = OpLabel
+ %42 = OpFunctionCall %void %textureDimensions_ba1481
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.glsl b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.glsl
index 3808a80..7387583 100644
--- a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.glsl
@@ -1,11 +1,170 @@
SKIP: FAILED
-../../src/tint/writer/glsl/generator_impl.cc:2544 internal compiler error: Multiplanar external texture transform was not run.
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+void textureLoad_8acf41() {
+ vec4 res = textureLoadExternal(arg_0_1, ext_tex_plane_1_1, ivec2(0, 0), ext_tex_params);
+}
+
+vec4 vertex_main() {
+ textureLoad_8acf41();
+ return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:36: 'textureLoadExternal' : no matching overloaded function found
+ERROR: 0:36: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:36: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+
+#version 310 es
+precision mediump float;
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+void textureLoad_8acf41() {
+ vec4 res = textureLoadExternal(arg_0_1, ext_tex_plane_1_1, ivec2(0, 0), ext_tex_params);
+}
+
+void fragment_main() {
+ textureLoad_8acf41();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:37: 'textureLoadExternal' : no matching overloaded function found
+ERROR: 0:37: '=' : cannot convert from ' const float' to ' temp mediump 4-component vector of float'
+ERROR: 0:37: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_1;
+void textureLoad_8acf41() {
+ vec4 res = textureLoadExternal(arg_0_1, ext_tex_plane_1_1, ivec2(0, 0), ext_tex_params);
+}
+
+void compute_main() {
+ textureLoad_8acf41();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:36: 'textureLoadExternal' : no matching overloaded function found
+ERROR: 0:36: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:36: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl
index b3cea20..a14b74b 100644
--- a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl
@@ -1,11 +1,68 @@
-SKIP: FAILED
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
-C:\src\tint2\src\writer\hlsl\generator_impl.cc:3632 internal compiler error: Multiplanar external texture transform was not run.
+Texture2D<float4> ext_tex_plane_1 : register(t1, space1);
+cbuffer cbuffer_ext_tex_params : register(b2, space1) {
+ uint4 ext_tex_params[2];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, int2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.Load(int3(coord, 0));
+ }
+ const float y = (plane0.Load(int3(coord, 0)).r - 0.0625f);
+ const float2 uv = (plane1.Load(int3(coord, 0)).rg - 0.5f);
+ const float u = uv.x;
+ const float v = uv.y;
+ const float r = ((1.164000034f * y) + (params.vr * v));
+ const float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ const float b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+ExternalTextureParams tint_symbol_1(uint4 buffer[2], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 4u)) / 4;
+ const uint scalar_offset_2 = ((offset + 8u)) / 4;
+ const uint scalar_offset_3 = ((offset + 12u)) / 4;
+ const uint scalar_offset_4 = ((offset + 16u)) / 4;
+ const ExternalTextureParams tint_symbol_4 = {buffer[scalar_offset / 4][scalar_offset % 4], asfloat(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4]), asfloat(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4]), asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4])};
+ return tint_symbol_4;
+}
+
+void textureLoad_8acf41() {
+ float4 res = textureLoadExternal(arg_0, ext_tex_plane_1, int2(0, 0), tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ textureLoad_8acf41();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ textureLoad_8acf41();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ textureLoad_8acf41();
+ return;
+}
diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.msl b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.msl
index 21d485b..d91b328 100644
--- a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.msl
+++ b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.msl
@@ -1,11 +1,55 @@
-SKIP: FAILED
+#include <metal_stdlib>
-C:\src\tint2\src\writer\msl\generator_impl.cc:2344 internal compiler error: Multiplanar external texture transform was not run.
+using namespace metal;
+struct ExternalTextureParams {
+ /* 0x0000 */ uint numPlanes;
+ /* 0x0004 */ float vr;
+ /* 0x0008 */ float ug;
+ /* 0x000c */ float vg;
+ /* 0x0010 */ float ub;
+};
+float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, int2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.read(uint2(coord), 0);
+ }
+ float const y = (plane0.read(uint2(coord), 0)[0] - 0.0625f);
+ float2 const uv = (float4(plane1.read(uint2(coord), 0)).rg - 0.5f);
+ float const u = uv[0];
+ float const v = uv[1];
+ float const r = ((1.164000034f * y) + (params.vr * v));
+ float const g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float const b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void textureLoad_8acf41(texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, const constant ExternalTextureParams* const tint_symbol_3) {
+ float4 res = textureLoadExternal(tint_symbol_1, tint_symbol_2, int2(), *(tint_symbol_3));
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_4, texture2d<float, access::sample> tint_symbol_5, const constant ExternalTextureParams* const tint_symbol_6) {
+ textureLoad_8acf41(tint_symbol_4, tint_symbol_5, tint_symbol_6);
+ return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_7 [[texture(0)]], texture2d<float, access::sample> tint_symbol_8 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_9 [[buffer(2)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8, tint_symbol_9);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_10 [[texture(0)]], texture2d<float, access::sample> tint_symbol_11 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_12 [[buffer(2)]]) {
+ textureLoad_8acf41(tint_symbol_10, tint_symbol_11, tint_symbol_12);
+ return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_13 [[texture(0)]], texture2d<float, access::sample> tint_symbol_14 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_15 [[buffer(2)]]) {
+ textureLoad_8acf41(tint_symbol_13, tint_symbol_14, tint_symbol_15);
+ return;
+}
+
diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm
index 87554bb..4ac3e99 100644
--- a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm
@@ -1,11 +1,157 @@
-SKIP: FAILED
-
-C:\src\tint2\src\writer\spirv\builder.cc:4013 internal compiler error: Multiplanar external texture transform was not run.
-
-
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 91
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %ext_tex_plane_1 "ext_tex_plane_1"
+ OpName %ExternalTextureParams "ExternalTextureParams"
+ OpMemberName %ExternalTextureParams 0 "numPlanes"
+ OpMemberName %ExternalTextureParams 1 "vr"
+ OpMemberName %ExternalTextureParams 2 "ug"
+ OpMemberName %ExternalTextureParams 3 "vg"
+ OpMemberName %ExternalTextureParams 4 "ub"
+ OpName %ext_tex_params "ext_tex_params"
+ OpName %arg_0 "arg_0"
+ OpName %textureLoadExternal "textureLoadExternal"
+ OpName %plane0 "plane0"
+ OpName %plane1 "plane1"
+ OpName %coord "coord"
+ OpName %params "params"
+ OpName %textureLoad_8acf41 "textureLoad_8acf41"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %ext_tex_plane_1 DescriptorSet 1
+ OpDecorate %ext_tex_plane_1 Binding 1
+ OpDecorate %ExternalTextureParams Block
+ OpMemberDecorate %ExternalTextureParams 0 Offset 0
+ OpMemberDecorate %ExternalTextureParams 1 Offset 4
+ OpMemberDecorate %ExternalTextureParams 2 Offset 8
+ OpMemberDecorate %ExternalTextureParams 3 Offset 12
+ OpMemberDecorate %ExternalTextureParams 4 Offset 16
+ OpDecorate %ext_tex_params NonWritable
+ OpDecorate %ext_tex_params DescriptorSet 1
+ OpDecorate %ext_tex_params Binding 2
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %uint = OpTypeInt 32 0
+%ExternalTextureParams = OpTypeStruct %uint %float %float %float %float
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%ext_tex_params = OpVariable %_ptr_Uniform_ExternalTextureParams Uniform
+ %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %17 = OpTypeFunction %v4float %11 %11 %v2int %ExternalTextureParams
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %int_0 = OpConstant %int 0
+%float_0_0625 = OpConstant %float 0.0625
+ %v2float = OpTypeVector %float 2
+ %float_0_5 = OpConstant %float 0.5
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+ %45 = OpConstantNull %v2float
+%float_1_16400003 = OpConstant %float 1.16400003
+ %float_1 = OpConstant %float 1
+ %void = OpTypeVoid
+ %67 = OpTypeFunction %void
+ %74 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+ %78 = OpTypeFunction %v4float
+%textureLoadExternal = OpFunction %v4float None %17
+ %plane0 = OpFunctionParameter %11
+ %plane1 = OpFunctionParameter %11
+ %coord = OpFunctionParameter %v2int
+ %params = OpFunctionParameter %ExternalTextureParams
+ %25 = OpLabel
+ %43 = OpVariable %_ptr_Function_v2float Function %45
+ %26 = OpCompositeExtract %uint %params 0
+ %28 = OpIEqual %bool %26 %uint_1
+ OpSelectionMerge %30 None
+ OpBranchConditional %28 %31 %30
+ %31 = OpLabel
+ %32 = OpImageFetch %v4float %plane0 %coord Lod %int_0
+ OpReturnValue %32
+ %30 = OpLabel
+ %34 = OpImageFetch %v4float %plane0 %coord Lod %int_0
+ %35 = OpCompositeExtract %float %34 0
+ %37 = OpFSub %float %35 %float_0_0625
+ %38 = OpImageFetch %v4float %plane1 %coord Lod %int_0
+ %40 = OpVectorShuffle %v2float %38 %38 0 1
+ %46 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+ %42 = OpFSub %v2float %40 %46
+ %47 = OpCompositeExtract %float %42 0
+ %48 = OpCompositeExtract %float %42 1
+ %50 = OpFMul %float %float_1_16400003 %37
+ %51 = OpCompositeExtract %float %params 1
+ %52 = OpFMul %float %51 %48
+ %53 = OpFAdd %float %50 %52
+ %54 = OpFMul %float %float_1_16400003 %37
+ %55 = OpCompositeExtract %float %params 2
+ %56 = OpFMul %float %55 %47
+ %57 = OpFSub %float %54 %56
+ %58 = OpCompositeExtract %float %params 3
+ %59 = OpFMul %float %58 %48
+ %60 = OpFSub %float %57 %59
+ %61 = OpFMul %float %float_1_16400003 %37
+ %62 = OpCompositeExtract %float %params 4
+ %63 = OpFMul %float %62 %47
+ %64 = OpFAdd %float %61 %63
+ %66 = OpCompositeConstruct %v4float %53 %60 %64 %float_1
+ OpReturnValue %66
+ OpFunctionEnd
+%textureLoad_8acf41 = OpFunction %void None %67
+ %70 = OpLabel
+ %res = OpVariable %_ptr_Function_v4float Function %5
+ %72 = OpLoad %11 %arg_0
+ %73 = OpLoad %11 %ext_tex_plane_1
+ %75 = OpLoad %ExternalTextureParams %ext_tex_params
+ %71 = OpFunctionCall %v4float %textureLoadExternal %72 %73 %74 %75
+ OpStore %res %71
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %78
+ %80 = OpLabel
+ %81 = OpFunctionCall %void %textureLoad_8acf41
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %67
+ %83 = OpLabel
+ %84 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %84
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %67
+ %86 = OpLabel
+ %87 = OpFunctionCall %void %textureLoad_8acf41
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %67
+ %89 = OpLabel
+ %90 = OpFunctionCall %void %textureLoad_8acf41
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.glsl b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.glsl
index 3808a80..3b57c67 100644
--- a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.glsl
+++ b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.glsl
@@ -1,11 +1,173 @@
SKIP: FAILED
-../../src/tint/writer/glsl/generator_impl.cc:2544 internal compiler error: Multiplanar external texture transform was not run.
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 3) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+vec4 textureSampleExternal(highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return textureLod(plane0_smp, coord, 0.0f);
+ }
+ float y = (textureLod(plane0_smp, coord, 0.0f).r - 0.0625f);
+ vec2 uv = (textureLod(plane1_smp, coord, 0.0f).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleLevel_979816() {
+ vec4 res = textureSampleExternal(arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f, 0.0f), ext_tex_params);
+}
+
+vec4 vertex_main() {
+ textureSampleLevel_979816();
+ return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:37: 'textureSampleExternal' : no matching overloaded function found
+ERROR: 0:37: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:37: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
+#version 310 es
+precision mediump float;
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 3) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return textureLod(plane0_smp, coord, 0.0f);
+ }
+ float y = (textureLod(plane0_smp, coord, 0.0f).r - 0.0625f);
+ vec2 uv = (textureLod(plane1_smp, coord, 0.0f).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleLevel_979816() {
+ vec4 res = textureSampleExternal(arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f, 0.0f), ext_tex_params);
+}
+
+void fragment_main() {
+ textureSampleLevel_979816();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:38: 'textureSampleExternal' : no matching overloaded function found
+ERROR: 0:38: '=' : cannot convert from ' const float' to ' temp mediump 4-component vector of float'
+ERROR: 0:38: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 3) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+
+vec4 textureSampleExternal(highp sampler2D plane0_smp, highp sampler2D plane1_smp, vec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return textureLod(plane0_smp, coord, 0.0f);
+ }
+ float y = (textureLod(plane0_smp, coord, 0.0f).r - 0.0625f);
+ vec2 uv = (textureLod(plane1_smp, coord, 0.0f).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+uniform highp sampler2D arg_0_arg_1;
+uniform highp sampler2D ext_tex_plane_1_arg_1;
+void textureSampleLevel_979816() {
+ vec4 res = textureSampleExternal(arg_0_arg_1, ext_tex_plane_1_arg_1, vec2(0.0f, 0.0f), ext_tex_params);
+}
+
+void compute_main() {
+ textureSampleLevel_979816();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:37: 'textureSampleExternal' : no matching overloaded function found
+ERROR: 0:37: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:37: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl
index b3cea20..4fe76bc 100644
--- a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl
+++ b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl
@@ -1,11 +1,69 @@
-SKIP: FAILED
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
-C:\src\tint2\src\writer\hlsl\generator_impl.cc:3632 internal compiler error: Multiplanar external texture transform was not run.
+Texture2D<float4> ext_tex_plane_1 : register(t2, space1);
+cbuffer cbuffer_ext_tex_params : register(b3, space1) {
+ uint4 ext_tex_params[2];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+SamplerState arg_1 : register(s1, space1);
+float4 textureSampleExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, SamplerState smp, float2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.SampleLevel(smp, coord, 0.0f);
+ }
+ const float y = (plane0.SampleLevel(smp, coord, 0.0f).r - 0.0625f);
+ const float2 uv = (plane1.SampleLevel(smp, coord, 0.0f).rg - 0.5f);
+ const float u = uv.x;
+ const float v = uv.y;
+ const float r = ((1.164000034f * y) + (params.vr * v));
+ const float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ const float b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+ExternalTextureParams tint_symbol_1(uint4 buffer[2], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 4u)) / 4;
+ const uint scalar_offset_2 = ((offset + 8u)) / 4;
+ const uint scalar_offset_3 = ((offset + 12u)) / 4;
+ const uint scalar_offset_4 = ((offset + 16u)) / 4;
+ const ExternalTextureParams tint_symbol_4 = {buffer[scalar_offset / 4][scalar_offset % 4], asfloat(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4]), asfloat(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4]), asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4])};
+ return tint_symbol_4;
+}
+
+void textureSampleLevel_979816() {
+ float4 res = textureSampleExternal(arg_0, ext_tex_plane_1, arg_1, float2(0.0f, 0.0f), tint_symbol_1(ext_tex_params, 0u));
+}
+
+struct tint_symbol {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ textureSampleLevel_979816();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol wrapper_result = (tint_symbol)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ textureSampleLevel_979816();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ textureSampleLevel_979816();
+ return;
+}
diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.msl b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.msl
index 21d485b..9b4be13 100644
--- a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.msl
+++ b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.msl
@@ -1,11 +1,55 @@
-SKIP: FAILED
+#include <metal_stdlib>
-C:\src\tint2\src\writer\msl\generator_impl.cc:2344 internal compiler error: Multiplanar external texture transform was not run.
+using namespace metal;
+struct ExternalTextureParams {
+ /* 0x0000 */ uint numPlanes;
+ /* 0x0004 */ float vr;
+ /* 0x0008 */ float ug;
+ /* 0x000c */ float vg;
+ /* 0x0010 */ float ub;
+};
+float4 textureSampleExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, sampler smp, float2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.sample(smp, coord, level(0.0f));
+ }
+ float const y = (plane0.sample(smp, coord, level(0.0f))[0] - 0.0625f);
+ float2 const uv = (float4(plane1.sample(smp, coord, level(0.0f))).rg - 0.5f);
+ float const u = uv[0];
+ float const v = uv[1];
+ float const r = ((1.164000034f * y) + (params.vr * v));
+ float const g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float const b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+void textureSampleLevel_979816(texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3, const constant ExternalTextureParams* const tint_symbol_4) {
+ float4 res = textureSampleExternal(tint_symbol_1, tint_symbol_2, tint_symbol_3, float2(), *(tint_symbol_4));
+}
+
+struct tint_symbol {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_5, texture2d<float, access::sample> tint_symbol_6, sampler tint_symbol_7, const constant ExternalTextureParams* const tint_symbol_8) {
+ textureSampleLevel_979816(tint_symbol_5, tint_symbol_6, tint_symbol_7, tint_symbol_8);
+ return float4();
+}
+
+vertex tint_symbol vertex_main(texture2d<float, access::sample> tint_symbol_9 [[texture(0)]], texture2d<float, access::sample> tint_symbol_10 [[texture(1)]], sampler tint_symbol_11 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_12 [[buffer(2)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_9, tint_symbol_10, tint_symbol_11, tint_symbol_12);
+ tint_symbol wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_13 [[texture(0)]], texture2d<float, access::sample> tint_symbol_14 [[texture(1)]], sampler tint_symbol_15 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_16 [[buffer(2)]]) {
+ textureSampleLevel_979816(tint_symbol_13, tint_symbol_14, tint_symbol_15, tint_symbol_16);
+ return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_17 [[texture(0)]], texture2d<float, access::sample> tint_symbol_18 [[texture(1)]], sampler tint_symbol_19 [[sampler(0)]], const constant ExternalTextureParams* tint_symbol_20 [[buffer(2)]]) {
+ textureSampleLevel_979816(tint_symbol_17, tint_symbol_18, tint_symbol_19, tint_symbol_20);
+ return;
+}
+
diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm
index 87554bb..62347b9 100644
--- a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm
+++ b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm
@@ -1,11 +1,167 @@
-SKIP: FAILED
-
-C:\src\tint2\src\writer\spirv\builder.cc:4013 internal compiler error: Multiplanar external texture transform was not run.
-
-
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 97
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %ext_tex_plane_1 "ext_tex_plane_1"
+ OpName %ExternalTextureParams "ExternalTextureParams"
+ OpMemberName %ExternalTextureParams 0 "numPlanes"
+ OpMemberName %ExternalTextureParams 1 "vr"
+ OpMemberName %ExternalTextureParams 2 "ug"
+ OpMemberName %ExternalTextureParams 3 "vg"
+ OpMemberName %ExternalTextureParams 4 "ub"
+ OpName %ext_tex_params "ext_tex_params"
+ OpName %arg_0 "arg_0"
+ OpName %arg_1 "arg_1"
+ OpName %textureSampleExternal "textureSampleExternal"
+ OpName %plane0 "plane0"
+ OpName %plane1 "plane1"
+ OpName %smp "smp"
+ OpName %coord "coord"
+ OpName %params "params"
+ OpName %textureSampleLevel_979816 "textureSampleLevel_979816"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %ext_tex_plane_1 DescriptorSet 1
+ OpDecorate %ext_tex_plane_1 Binding 2
+ OpDecorate %ExternalTextureParams Block
+ OpMemberDecorate %ExternalTextureParams 0 Offset 0
+ OpMemberDecorate %ExternalTextureParams 1 Offset 4
+ OpMemberDecorate %ExternalTextureParams 2 Offset 8
+ OpMemberDecorate %ExternalTextureParams 3 Offset 12
+ OpMemberDecorate %ExternalTextureParams 4 Offset 16
+ OpDecorate %ext_tex_params NonWritable
+ OpDecorate %ext_tex_params DescriptorSet 1
+ OpDecorate %ext_tex_params Binding 3
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ OpDecorate %arg_1 DescriptorSet 1
+ OpDecorate %arg_1 Binding 1
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %uint = OpTypeInt 32 0
+%ExternalTextureParams = OpTypeStruct %uint %float %float %float %float
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%ext_tex_params = OpVariable %_ptr_Uniform_ExternalTextureParams Uniform
+ %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %19 = OpTypeSampler
+%_ptr_UniformConstant_19 = OpTypePointer UniformConstant %19
+ %arg_1 = OpVariable %_ptr_UniformConstant_19 UniformConstant
+ %v2float = OpTypeVector %float 2
+ %20 = OpTypeFunction %v4float %11 %11 %19 %v2float %ExternalTextureParams
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %36 = OpTypeSampledImage %11
+ %float_0 = OpConstant %float 0
+%float_0_0625 = OpConstant %float 0.0625
+ %float_0_5 = OpConstant %float 0.5
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+ %51 = OpConstantNull %v2float
+%float_1_16400003 = OpConstant %float 1.16400003
+ %float_1 = OpConstant %float 1
+ %void = OpTypeVoid
+ %73 = OpTypeFunction %void
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+ %84 = OpTypeFunction %v4float
+%textureSampleExternal = OpFunction %v4float None %20
+ %plane0 = OpFunctionParameter %11
+ %plane1 = OpFunctionParameter %11
+ %smp = OpFunctionParameter %19
+ %coord = OpFunctionParameter %v2float
+ %params = OpFunctionParameter %ExternalTextureParams
+ %28 = OpLabel
+ %49 = OpVariable %_ptr_Function_v2float Function %51
+ %29 = OpCompositeExtract %uint %params 0
+ %31 = OpIEqual %bool %29 %uint_1
+ OpSelectionMerge %33 None
+ OpBranchConditional %31 %34 %33
+ %34 = OpLabel
+ %37 = OpSampledImage %36 %plane0 %smp
+ %35 = OpImageSampleExplicitLod %v4float %37 %coord Lod %float_0
+ OpReturnValue %35
+ %33 = OpLabel
+ %40 = OpSampledImage %36 %plane0 %smp
+ %39 = OpImageSampleExplicitLod %v4float %40 %coord Lod %float_0
+ %41 = OpCompositeExtract %float %39 0
+ %43 = OpFSub %float %41 %float_0_0625
+ %45 = OpSampledImage %36 %plane1 %smp
+ %44 = OpImageSampleExplicitLod %v4float %45 %coord Lod %float_0
+ %46 = OpVectorShuffle %v2float %44 %44 0 1
+ %52 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+ %48 = OpFSub %v2float %46 %52
+ %53 = OpCompositeExtract %float %48 0
+ %54 = OpCompositeExtract %float %48 1
+ %56 = OpFMul %float %float_1_16400003 %43
+ %57 = OpCompositeExtract %float %params 1
+ %58 = OpFMul %float %57 %54
+ %59 = OpFAdd %float %56 %58
+ %60 = OpFMul %float %float_1_16400003 %43
+ %61 = OpCompositeExtract %float %params 2
+ %62 = OpFMul %float %61 %53
+ %63 = OpFSub %float %60 %62
+ %64 = OpCompositeExtract %float %params 3
+ %65 = OpFMul %float %64 %54
+ %66 = OpFSub %float %63 %65
+ %67 = OpFMul %float %float_1_16400003 %43
+ %68 = OpCompositeExtract %float %params 4
+ %69 = OpFMul %float %68 %53
+ %70 = OpFAdd %float %67 %69
+ %72 = OpCompositeConstruct %v4float %59 %66 %70 %float_1
+ OpReturnValue %72
+ OpFunctionEnd
+%textureSampleLevel_979816 = OpFunction %void None %73
+ %76 = OpLabel
+ %res = OpVariable %_ptr_Function_v4float Function %5
+ %78 = OpLoad %11 %arg_0
+ %79 = OpLoad %11 %ext_tex_plane_1
+ %80 = OpLoad %19 %arg_1
+ %81 = OpLoad %ExternalTextureParams %ext_tex_params
+ %77 = OpFunctionCall %v4float %textureSampleExternal %78 %79 %80 %51 %81
+ OpStore %res %77
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %84
+ %86 = OpLabel
+ %87 = OpFunctionCall %void %textureSampleLevel_979816
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %73
+ %89 = OpLabel
+ %90 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %90
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %73
+ %92 = OpLabel
+ %93 = OpFunctionCall %void %textureSampleLevel_979816
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %73
+ %95 = OpLabel
+ %96 = OpFunctionCall %void %textureSampleLevel_979816
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl b/test/tint/builtins/textureLoad/texture_external_param.wgsl
new file mode 100644
index 0000000..4a8edf8
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl
@@ -0,0 +1,25 @@
+@group(1) @binding(0) var arg_0: texture_external;
+
+fn textureLoad2d(texture: texture_external, coords: vec2<i32>) -> vec4<f32> {
+ return textureLoad(texture, coords);
+}
+
+fn doTextureLoad() {
+ var res: vec4<f32> = textureLoad2d(arg_0, vec2<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ doTextureLoad();
+ return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+ doTextureLoad();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+ doTextureLoad();
+}
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.glsl b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.glsl
new file mode 100644
index 0000000..17d620f
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.glsl
@@ -0,0 +1,182 @@
+SKIP: FAILED
+
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+vec4 textureLoad2d(highp sampler2D tint_symbol_1, highp sampler2D ext_tex_plane_1_1_1, ExternalTextureParams ext_tex_params_1, ivec2 coords) {
+ return textureLoadExternal(tint_symbol_1, ext_tex_plane_1_1_1, coords, ext_tex_params_1);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_2;
+void doTextureLoad() {
+ vec4 res = textureLoad2d(arg_0_1, ext_tex_plane_1_2, ext_tex_params, ivec2(0, 0));
+}
+
+vec4 vertex_main() {
+ doTextureLoad();
+ return vec4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+void main() {
+ vec4 inner_result = vertex_main();
+ gl_Position = inner_result;
+ gl_Position.y = -(gl_Position.y);
+ gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:40: 'textureLoad2d' : no matching overloaded function found
+ERROR: 0:40: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:40: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
+#version 310 es
+precision mediump float;
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+vec4 textureLoad2d(highp sampler2D tint_symbol_1, highp sampler2D ext_tex_plane_1_1_1, ExternalTextureParams ext_tex_params_1, ivec2 coords) {
+ return textureLoadExternal(tint_symbol_1, ext_tex_plane_1_1_1, coords, ext_tex_params_1);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_2;
+void doTextureLoad() {
+ vec4 res = textureLoad2d(arg_0_1, ext_tex_plane_1_2, ext_tex_params, ivec2(0, 0));
+}
+
+void fragment_main() {
+ doTextureLoad();
+}
+
+void main() {
+ fragment_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:41: 'textureLoad2d' : no matching overloaded function found
+ERROR: 0:41: '=' : cannot convert from ' const float' to ' temp mediump 4-component vector of float'
+ERROR: 0:41: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
+#version 310 es
+
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+layout(binding = 2) uniform ExternalTextureParams_1 {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+} ext_tex_params;
+
+vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return texelFetch(plane0_1, coord, 0);
+ }
+ float y = (texelFetch(plane0_1, coord, 0).r - 0.0625f);
+ vec2 uv = (texelFetch(plane1_1, coord, 0).rg - 0.5f);
+ float u = uv.x;
+ float v = uv.y;
+ float r = ((1.164000034f * y) + (params.vr * v));
+ float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float b = ((1.164000034f * y) + (params.ub * u));
+ return vec4(r, g, b, 1.0f);
+}
+
+vec4 textureLoad2d(highp sampler2D tint_symbol_1, highp sampler2D ext_tex_plane_1_1_1, ExternalTextureParams ext_tex_params_1, ivec2 coords) {
+ return textureLoadExternal(tint_symbol_1, ext_tex_plane_1_1_1, coords, ext_tex_params_1);
+}
+
+uniform highp sampler2D arg_0_1;
+uniform highp sampler2D ext_tex_plane_1_2;
+void doTextureLoad() {
+ vec4 res = textureLoad2d(arg_0_1, ext_tex_plane_1_2, ext_tex_params, ivec2(0, 0));
+}
+
+void compute_main() {
+ doTextureLoad();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main();
+ return;
+}
+Error parsing GLSL shader:
+ERROR: 0:40: 'textureLoad2d' : no matching overloaded function found
+ERROR: 0:40: '=' : cannot convert from ' const float' to ' temp highp 4-component vector of float'
+ERROR: 0:40: '' : compilation terminated
+ERROR: 3 compilation errors. No code generated.
+
+
+
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.hlsl b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.hlsl
new file mode 100644
index 0000000..e7f6715
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.hlsl
@@ -0,0 +1,72 @@
+struct ExternalTextureParams {
+ uint numPlanes;
+ float vr;
+ float ug;
+ float vg;
+ float ub;
+};
+
+Texture2D<float4> ext_tex_plane_1 : register(t1, space1);
+cbuffer cbuffer_ext_tex_params : register(b2, space1) {
+ uint4 ext_tex_params[2];
+};
+Texture2D<float4> arg_0 : register(t0, space1);
+
+float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, int2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.Load(int3(coord, 0));
+ }
+ const float y = (plane0.Load(int3(coord, 0)).r - 0.0625f);
+ const float2 uv = (plane1.Load(int3(coord, 0)).rg - 0.5f);
+ const float u = uv.x;
+ const float v = uv.y;
+ const float r = ((1.164000034f * y) + (params.vr * v));
+ const float g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ const float b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
+
+float4 textureLoad2d(Texture2D<float4> tint_symbol, Texture2D<float4> ext_tex_plane_1_1, ExternalTextureParams ext_tex_params_1, int2 coords) {
+ return textureLoadExternal(tint_symbol, ext_tex_plane_1_1, coords, ext_tex_params_1);
+}
+
+ExternalTextureParams tint_symbol_2(uint4 buffer[2], uint offset) {
+ const uint scalar_offset = ((offset + 0u)) / 4;
+ const uint scalar_offset_1 = ((offset + 4u)) / 4;
+ const uint scalar_offset_2 = ((offset + 8u)) / 4;
+ const uint scalar_offset_3 = ((offset + 12u)) / 4;
+ const uint scalar_offset_4 = ((offset + 16u)) / 4;
+ const ExternalTextureParams tint_symbol_5 = {buffer[scalar_offset / 4][scalar_offset % 4], asfloat(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4]), asfloat(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4]), asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4])};
+ return tint_symbol_5;
+}
+
+void doTextureLoad() {
+ float4 res = textureLoad2d(arg_0, ext_tex_plane_1, tint_symbol_2(ext_tex_params, 0u), int2(0, 0));
+}
+
+struct tint_symbol_1 {
+ float4 value : SV_Position;
+};
+
+float4 vertex_main_inner() {
+ doTextureLoad();
+ return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol_1 vertex_main() {
+ const float4 inner_result = vertex_main_inner();
+ tint_symbol_1 wrapper_result = (tint_symbol_1)0;
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+void fragment_main() {
+ doTextureLoad();
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ doTextureLoad();
+ return;
+}
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.msl b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.msl
new file mode 100644
index 0000000..4e82e4d
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.msl
@@ -0,0 +1,59 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct ExternalTextureParams {
+ /* 0x0000 */ uint numPlanes;
+ /* 0x0004 */ float vr;
+ /* 0x0008 */ float ug;
+ /* 0x000c */ float vg;
+ /* 0x0010 */ float ub;
+};
+
+float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, int2 coord, ExternalTextureParams params) {
+ if ((params.numPlanes == 1u)) {
+ return plane0.read(uint2(coord), 0);
+ }
+ float const y = (plane0.read(uint2(coord), 0)[0] - 0.0625f);
+ float2 const uv = (float4(plane1.read(uint2(coord), 0)).rg - 0.5f);
+ float const u = uv[0];
+ float const v = uv[1];
+ float const r = ((1.164000034f * y) + (params.vr * v));
+ float const g = (((1.164000034f * y) - (params.ug * u)) - (params.vg * v));
+ float const b = ((1.164000034f * y) + (params.ub * u));
+ return float4(r, g, b, 1.0f);
+}
+
+float4 textureLoad2d(texture2d<float, access::sample> tint_symbol, texture2d<float, access::sample> ext_tex_plane_1_1, ExternalTextureParams ext_tex_params_1, int2 coords) {
+ return textureLoadExternal(tint_symbol, ext_tex_plane_1_1, coords, ext_tex_params_1);
+}
+
+void doTextureLoad(texture2d<float, access::sample> tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, const constant ExternalTextureParams* const tint_symbol_4) {
+ float4 res = textureLoad2d(tint_symbol_2, tint_symbol_3, *(tint_symbol_4), int2());
+}
+
+struct tint_symbol_1 {
+ float4 value [[position]];
+};
+
+float4 vertex_main_inner(texture2d<float, access::sample> tint_symbol_5, texture2d<float, access::sample> tint_symbol_6, const constant ExternalTextureParams* const tint_symbol_7) {
+ doTextureLoad(tint_symbol_5, tint_symbol_6, tint_symbol_7);
+ return float4();
+}
+
+vertex tint_symbol_1 vertex_main(texture2d<float, access::sample> tint_symbol_8 [[texture(0)]], texture2d<float, access::sample> tint_symbol_9 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_10 [[buffer(2)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_8, tint_symbol_9, tint_symbol_10);
+ tint_symbol_1 wrapper_result = {};
+ wrapper_result.value = inner_result;
+ return wrapper_result;
+}
+
+fragment void fragment_main(texture2d<float, access::sample> tint_symbol_11 [[texture(0)]], texture2d<float, access::sample> tint_symbol_12 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_13 [[buffer(2)]]) {
+ doTextureLoad(tint_symbol_11, tint_symbol_12, tint_symbol_13);
+ return;
+}
+
+kernel void compute_main(texture2d<float, access::sample> tint_symbol_14 [[texture(0)]], texture2d<float, access::sample> tint_symbol_15 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_16 [[buffer(2)]]) {
+ doTextureLoad(tint_symbol_14, tint_symbol_15, tint_symbol_16);
+ return;
+}
+
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.spvasm b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.spvasm
new file mode 100644
index 0000000..cd8297f
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.spvasm
@@ -0,0 +1,172 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 99
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %value "value"
+ OpName %vertex_point_size "vertex_point_size"
+ OpName %ext_tex_plane_1 "ext_tex_plane_1"
+ OpName %ExternalTextureParams "ExternalTextureParams"
+ OpMemberName %ExternalTextureParams 0 "numPlanes"
+ OpMemberName %ExternalTextureParams 1 "vr"
+ OpMemberName %ExternalTextureParams 2 "ug"
+ OpMemberName %ExternalTextureParams 3 "vg"
+ OpMemberName %ExternalTextureParams 4 "ub"
+ OpName %ext_tex_params "ext_tex_params"
+ OpName %arg_0 "arg_0"
+ OpName %textureLoadExternal "textureLoadExternal"
+ OpName %plane0 "plane0"
+ OpName %plane1 "plane1"
+ OpName %coord "coord"
+ OpName %params "params"
+ OpName %textureLoad2d "textureLoad2d"
+ OpName %texture "texture"
+ OpName %ext_tex_plane_1_1 "ext_tex_plane_1_1"
+ OpName %ext_tex_params_1 "ext_tex_params_1"
+ OpName %coords "coords"
+ OpName %doTextureLoad "doTextureLoad"
+ OpName %res "res"
+ OpName %vertex_main_inner "vertex_main_inner"
+ OpName %vertex_main "vertex_main"
+ OpName %fragment_main "fragment_main"
+ OpName %compute_main "compute_main"
+ OpDecorate %value BuiltIn Position
+ OpDecorate %vertex_point_size BuiltIn PointSize
+ OpDecorate %ext_tex_plane_1 DescriptorSet 1
+ OpDecorate %ext_tex_plane_1 Binding 1
+ OpDecorate %ExternalTextureParams Block
+ OpMemberDecorate %ExternalTextureParams 0 Offset 0
+ OpMemberDecorate %ExternalTextureParams 1 Offset 4
+ OpMemberDecorate %ExternalTextureParams 2 Offset 8
+ OpMemberDecorate %ExternalTextureParams 3 Offset 12
+ OpMemberDecorate %ExternalTextureParams 4 Offset 16
+ OpDecorate %ext_tex_params NonWritable
+ OpDecorate %ext_tex_params DescriptorSet 1
+ OpDecorate %ext_tex_params Binding 2
+ OpDecorate %arg_0 DescriptorSet 1
+ OpDecorate %arg_0 Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+ %5 = OpConstantNull %v4float
+ %value = OpVariable %_ptr_Output_v4float Output %5
+%_ptr_Output_float = OpTypePointer Output %float
+ %8 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %8
+ %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %uint = OpTypeInt 32 0
+%ExternalTextureParams = OpTypeStruct %uint %float %float %float %float
+%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams
+%ext_tex_params = OpVariable %_ptr_Uniform_ExternalTextureParams Uniform
+ %arg_0 = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %int = OpTypeInt 32 1
+ %v2int = OpTypeVector %int 2
+ %17 = OpTypeFunction %v4float %11 %11 %v2int %ExternalTextureParams
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %int_0 = OpConstant %int 0
+%float_0_0625 = OpConstant %float 0.0625
+ %v2float = OpTypeVector %float 2
+ %float_0_5 = OpConstant %float 0.5
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+ %45 = OpConstantNull %v2float
+%float_1_16400003 = OpConstant %float 1.16400003
+ %float_1 = OpConstant %float 1
+ %67 = OpTypeFunction %v4float %11 %11 %ExternalTextureParams %v2int
+ %void = OpTypeVoid
+ %75 = OpTypeFunction %void
+ %83 = OpConstantNull %v2int
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+ %86 = OpTypeFunction %v4float
+%textureLoadExternal = OpFunction %v4float None %17
+ %plane0 = OpFunctionParameter %11
+ %plane1 = OpFunctionParameter %11
+ %coord = OpFunctionParameter %v2int
+ %params = OpFunctionParameter %ExternalTextureParams
+ %25 = OpLabel
+ %43 = OpVariable %_ptr_Function_v2float Function %45
+ %26 = OpCompositeExtract %uint %params 0
+ %28 = OpIEqual %bool %26 %uint_1
+ OpSelectionMerge %30 None
+ OpBranchConditional %28 %31 %30
+ %31 = OpLabel
+ %32 = OpImageFetch %v4float %plane0 %coord Lod %int_0
+ OpReturnValue %32
+ %30 = OpLabel
+ %34 = OpImageFetch %v4float %plane0 %coord Lod %int_0
+ %35 = OpCompositeExtract %float %34 0
+ %37 = OpFSub %float %35 %float_0_0625
+ %38 = OpImageFetch %v4float %plane1 %coord Lod %int_0
+ %40 = OpVectorShuffle %v2float %38 %38 0 1
+ %46 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+ %42 = OpFSub %v2float %40 %46
+ %47 = OpCompositeExtract %float %42 0
+ %48 = OpCompositeExtract %float %42 1
+ %50 = OpFMul %float %float_1_16400003 %37
+ %51 = OpCompositeExtract %float %params 1
+ %52 = OpFMul %float %51 %48
+ %53 = OpFAdd %float %50 %52
+ %54 = OpFMul %float %float_1_16400003 %37
+ %55 = OpCompositeExtract %float %params 2
+ %56 = OpFMul %float %55 %47
+ %57 = OpFSub %float %54 %56
+ %58 = OpCompositeExtract %float %params 3
+ %59 = OpFMul %float %58 %48
+ %60 = OpFSub %float %57 %59
+ %61 = OpFMul %float %float_1_16400003 %37
+ %62 = OpCompositeExtract %float %params 4
+ %63 = OpFMul %float %62 %47
+ %64 = OpFAdd %float %61 %63
+ %66 = OpCompositeConstruct %v4float %53 %60 %64 %float_1
+ OpReturnValue %66
+ OpFunctionEnd
+%textureLoad2d = OpFunction %v4float None %67
+ %texture = OpFunctionParameter %11
+%ext_tex_plane_1_1 = OpFunctionParameter %11
+%ext_tex_params_1 = OpFunctionParameter %ExternalTextureParams
+ %coords = OpFunctionParameter %v2int
+ %73 = OpLabel
+ %74 = OpFunctionCall %v4float %textureLoadExternal %texture %ext_tex_plane_1_1 %coords %ext_tex_params_1
+ OpReturnValue %74
+ OpFunctionEnd
+%doTextureLoad = OpFunction %void None %75
+ %78 = OpLabel
+ %res = OpVariable %_ptr_Function_v4float Function %5
+ %80 = OpLoad %11 %arg_0
+ %81 = OpLoad %11 %ext_tex_plane_1
+ %82 = OpLoad %ExternalTextureParams %ext_tex_params
+ %79 = OpFunctionCall %v4float %textureLoad2d %80 %81 %82 %83
+ OpStore %res %79
+ OpReturn
+ OpFunctionEnd
+%vertex_main_inner = OpFunction %v4float None %86
+ %88 = OpLabel
+ %89 = OpFunctionCall %void %doTextureLoad
+ OpReturnValue %5
+ OpFunctionEnd
+%vertex_main = OpFunction %void None %75
+ %91 = OpLabel
+ %92 = OpFunctionCall %v4float %vertex_main_inner
+ OpStore %value %92
+ OpStore %vertex_point_size %float_1
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %75
+ %94 = OpLabel
+ %95 = OpFunctionCall %void %doTextureLoad
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %75
+ %97 = OpLabel
+ %98 = OpFunctionCall %void %doTextureLoad
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.wgsl b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.wgsl
new file mode 100644
index 0000000..edba45d
--- /dev/null
+++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.wgsl
@@ -0,0 +1,25 @@
+@group(1) @binding(0) var arg_0 : texture_external;
+
+fn textureLoad2d(texture : texture_external, coords : vec2<i32>) -> vec4<f32> {
+ return textureLoad(texture, coords);
+}
+
+fn doTextureLoad() {
+ var res : vec4<f32> = textureLoad2d(arg_0, vec2<i32>());
+}
+
+@stage(vertex)
+fn vertex_main() -> @builtin(position) vec4<f32> {
+ doTextureLoad();
+ return vec4<f32>();
+}
+
+@stage(fragment)
+fn fragment_main() {
+ doTextureLoad();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+ doTextureLoad();
+}
