From a730eb738e9f00fb52e9ac38cebe978373602a1e Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Wed, 6 Apr 2022 13:57:54 +0000 Subject: [PATCH] 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 Reviewed-by: Ben Clayton Commit-Queue: Antonio Maiorano --- src/tint/BUILD.gn | 2 + src/tint/CMakeLists.txt | 3 + src/tint/cmd/main.cc | 4 + src/tint/fuzzers/data_builder.h | 1 + src/tint/transform/glsl.cc | 18 +- src/tint/transform/glsl.h | 8 +- .../generate_external_texture_bindings.cc | 59 ++++++ .../generate_external_texture_bindings.h | 27 +++ ...generate_external_texture_bindings_test.cc | 131 +++++++++++++ src/tint/writer/glsl/generator.cc | 4 +- src/tint/writer/glsl/generator.h | 3 + src/tint/writer/hlsl/generator.cc | 1 + src/tint/writer/hlsl/generator.h | 2 + src/tint/writer/hlsl/generator_impl.cc | 9 + src/tint/writer/hlsl/generator_impl.h | 1 + src/tint/writer/hlsl/test_helper.h | 8 +- src/tint/writer/msl/generator.cc | 1 + src/tint/writer/msl/generator.h | 3 + src/tint/writer/msl/generator_impl.cc | 10 + src/tint/writer/msl/generator_impl.h | 1 + src/tint/writer/msl/test_helper.h | 1 + src/tint/writer/spirv/builder.cc | 11 +- src/tint/writer/spirv/builder.h | 3 +- src/tint/writer/spirv/generator.cc | 3 +- src/tint/writer/spirv/generator.h | 3 + test/tint/BUILD.gn | 1 + .../ba1481.wgsl.expected.glsl | 104 +++++++++- .../ba1481.wgsl.expected.hlsl | 43 ++++- .../ba1481.wgsl.expected.msl | 46 ++++- .../ba1481.wgsl.expected.spvasm | 111 +++++++++-- .../gen/textureLoad/8acf41.wgsl.expected.glsl | 173 ++++++++++++++++- .../gen/textureLoad/8acf41.wgsl.expected.hlsl | 73 ++++++- .../gen/textureLoad/8acf41.wgsl.expected.msl | 60 +++++- .../textureLoad/8acf41.wgsl.expected.spvasm | 168 ++++++++++++++-- .../979816.wgsl.expected.glsl | 176 ++++++++++++++++- .../979816.wgsl.expected.hlsl | 74 ++++++- .../979816.wgsl.expected.msl | 60 +++++- .../979816.wgsl.expected.spvasm | 178 +++++++++++++++-- .../textureLoad/texture_external_param.wgsl | 25 +++ .../texture_external_param.wgsl.expected.glsl | 182 ++++++++++++++++++ .../texture_external_param.wgsl.expected.hlsl | 72 +++++++ .../texture_external_param.wgsl.expected.msl | 59 ++++++ ...exture_external_param.wgsl.expected.spvasm | 172 +++++++++++++++++ .../texture_external_param.wgsl.expected.wgsl | 25 +++ 44 files changed, 2006 insertions(+), 113 deletions(-) create mode 100644 src/tint/writer/generate_external_texture_bindings.cc create mode 100644 src/tint/writer/generate_external_texture_bindings.h create mode 100644 src/tint/writer/generate_external_texture_bindings_test.cc create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.glsl create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.hlsl create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.msl create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.spvasm create mode 100644 test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.wgsl diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 6c525c858a..58ec6a4285 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -527,6 +527,8 @@ libtint_source_set("libtint_core_all_src") { "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 207b988821..138dfde824 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -451,6 +451,8 @@ set(TINT_LIB_SRCS 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 @@ if(TINT_BUILD_TESTS) 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 ff3587c85c..8a0d2935d0 100644 --- a/src/tint/cmd/main.cc +++ b/src/tint/cmd/main.cc @@ -569,6 +569,7 @@ bool GenerateSpirv(const tint::Program* program, const Options& options) { // 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 @@ bool GenerateMsl(const tint::Program* program, const Options& options) { // 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 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) { // 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 @@ bool GenerateGlsl(const tint::Program* program, const Options& options) { 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 31a1a98047..e0c104a9c8 100644 --- a/src/tint/fuzzers/data_builder.h +++ b/src/tint/fuzzers/data_builder.h @@ -172,6 +172,7 @@ class DataBuilder { 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 5bb0006b83..df6b6b3b61 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 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const { manager.Add(); manager.Add(); + + if (cfg && cfg->generate_external_texture_bindings) { + auto new_bindings_map = writer::GenerateExternalTextureBindings(in); + data.Add(new_bindings_map); + } + manager.Add(); + manager.Add(); if (auto* binding_info = inputs.Get()) { data.Add(*binding_info); @@ -106,6 +114,7 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const { BindingRemapper::AccessControls ac; data.Add(bp, ac, /* mayCollide */ true); } + manager.Add(); manager.Add(); @@ -124,8 +133,13 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const { 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 baec29a5fc..414f871055 100644 --- a/src/tint/transform/glsl.h +++ b/src/tint/transform/glsl.h @@ -37,8 +37,11 @@ class Glsl final : public Castable { /// @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 @@ class Glsl final : public Castable { /// 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 0000000000..44bd262d74 --- /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 +#include +#include + +#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 group_to_next_binding_number; + std::vector ext_tex_bps; + for (auto* var : program->AST().GlobalVariables()) { + if (auto* sem_var = program->Sem().Get(var)->As()) { + 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()) { + 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 0000000000..6f321d5557 --- /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 0000000000..3163007043 --- /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 + +#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 ab7079fdea..9ae88adc1b 100644 --- a/src/tint/writer/glsl/generator.cc +++ b/src/tint/writer/glsl/generator.cc @@ -43,7 +43,9 @@ Result Generate(const Program* program, options.allow_collisions); data.Add( options.binding_map, options.placeholder_binding_point); - data.Add(entry_point); + data.Add(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 08b3e93767..88be42993b 100644 --- a/src/tint/writer/glsl/generator.h +++ b/src/tint/writer/glsl/generator.h @@ -71,6 +71,9 @@ struct Options { /// 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 536591fb98..b6aa19c433 100644 --- a/src/tint/writer/hlsl/generator.cc +++ b/src/tint/writer/hlsl/generator.cc @@ -35,6 +35,7 @@ Result Generate(const Program* program, const Options& options) { // 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 63df62e032..aa7cfa08a7 100644 --- a/src/tint/writer/hlsl/generator.h +++ b/src/tint/writer/hlsl/generator.h @@ -53,6 +53,8 @@ struct Options { 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 a8981f9ebb..fccbdcd473 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 @@ SanitizedResult Sanitize( 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 @@ SanitizedResult Sanitize( 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( + new_bindings_map); + } + manager.Add(); + manager.Add(); // LocalizeStructArrayAssignment must come after: diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h index e3bb1b3a4e..84b533d9c5 100644 --- a/src/tint/writer/hlsl/generator_impl.h +++ b/src/tint/writer/hlsl/generator_impl.h @@ -77,6 +77,7 @@ SanitizedResult Sanitize( 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 5ef28a267c..cac464506d 100644 --- a/src/tint/writer/hlsl/test_helper.h +++ b/src/tint/writer/hlsl/test_helper.h @@ -78,9 +78,11 @@ class TestHelperBase : public BODY, public ProgramBuilder { << 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 fa3c2f07c7..715af8d392 100644 --- a/src/tint/writer/msl/generator.cc +++ b/src/tint/writer/msl/generator.cc @@ -38,6 +38,7 @@ Result Generate(const Program* program, const Options& options) { 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 aa7f8c5857..27f99683d9 100644 --- a/src/tint/writer/msl/generator.h +++ b/src/tint/writer/msl/generator.h @@ -61,6 +61,9 @@ struct Options { /// 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 76d19ee7f9..5807b70eb3 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 @@ class ScopedBitCast { private: std::ostream& s; }; + } // namespace SanitizedResult::SanitizedResult() = default; @@ -125,6 +127,7 @@ SanitizedResult Sanitize( 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 @@ SanitizedResult Sanitize( transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, fixed_sample_mask, emit_vertex_point_size); + if (generate_external_texture_bindings) { + auto new_bindings_map = GenerateExternalTextureBindings(in); + data.Add( + new_bindings_map); + } + manager.Add(); + manager.Add(); if (!disable_workgroup_init) { diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h index 17ad9bd6b1..c973fe622f 100644 --- a/src/tint/writer/msl/generator_impl.h +++ b/src/tint/writer/msl/generator_impl.h @@ -84,6 +84,7 @@ SanitizedResult Sanitize( 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 e1d383dae2..2e0f3f1eb2 100644 --- a/src/tint/writer/msl/test_helper.h +++ b/src/tint/writer/msl/test_helper.h @@ -79,6 +79,7 @@ class TestHelperBase : public BASE, public ProgramBuilder { 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 52ba914a5f..5de6e46ce9 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 @@ const sem::Type* ElementTypeOf(const sem::Type* ty) { 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 @@ SanitizedResult Sanitize(const Program* in, manager.Add(); } + if (generate_external_texture_bindings) { + auto new_bindings_map = GenerateExternalTextureBindings(in); + data.Add( + new_bindings_map); + } + manager.Add(); + manager.Add(); if (!disable_workgroup_init) { manager.Add(); diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h index ef48afd262..2a94635c17 100644 --- a/src/tint/writer/spirv/builder.h +++ b/src/tint/writer/spirv/builder.h @@ -65,7 +65,8 @@ struct SanitizedResult { /// @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 4b3b72e454..f27514aa0e 100644 --- a/src/tint/writer/spirv/generator.cc +++ b/src/tint/writer/spirv/generator.cc @@ -32,7 +32,8 @@ Result Generate(const Program* program, const Options& options) { 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 7cf9654b1e..191e43abfa 100644 --- a/src/tint/writer/spirv/generator.h +++ b/src/tint/writer/spirv/generator.h @@ -42,6 +42,9 @@ struct Options { /// 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 b6bcdd49a8..d2ca28869c 100644 --- a/test/tint/BUILD.gn +++ b/test/tint/BUILD.gn @@ -376,6 +376,7 @@ tint_unittests_source_set("tint_unittests_writer_src") { 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 3808a80307..b351a3392e 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 b3cea20776..4697e5b716 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 ext_tex_plane_1 : register(t1, space1); +cbuffer cbuffer_ext_tex_params : register(b2, space1) { + uint4 ext_tex_params[2]; +}; +Texture2D 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 21d485b7ec..1faa96afe7 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 -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 tint_symbol_1) { + int2 res = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height()); +} + +struct tint_symbol { + float4 value [[position]]; +}; + +float4 vertex_main_inner(texture2d tint_symbol_2) { + textureDimensions_ba1481(tint_symbol_2); + return float4(); +} + +vertex tint_symbol vertex_main(texture2d 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 tint_symbol_4 [[texture(0)]]) { + textureDimensions_ba1481(tint_symbol_4); + return; +} + +kernel void compute_main(texture2d tint_symbol_5 [[texture(0)]]) { + textureDimensions_ba1481(tint_symbol_5); + return; +} -******************************************************************** -* 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. * -******************************************************************** diff --git a/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm b/test/tint/builtins/gen/textureDimensions/ba1481.wgsl.expected.spvasm index 87554bb606..461515fd09 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 3808a80307..7387583915 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. + + + +#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. + -******************************************************************** -* 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. * -******************************************************************** diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.hlsl index b3cea20776..a14b74b1db 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 ext_tex_plane_1 : register(t1, space1); +cbuffer cbuffer_ext_tex_params : register(b2, space1) { + uint4 ext_tex_params[2]; +}; +Texture2D arg_0 : register(t0, space1); +float4 textureLoadExternal(Texture2D plane0, Texture2D 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 21d485b7ec..d91b328b05 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 -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 plane0, texture2d 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); +} + +void textureLoad_8acf41(texture2d tint_symbol_1, texture2d 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 tint_symbol_4, texture2d 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 tint_symbol_7 [[texture(0)]], texture2d 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 tint_symbol_10 [[texture(0)]], texture2d 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 tint_symbol_13 [[texture(0)]], texture2d 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; +} -******************************************************************** -* 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. * -******************************************************************** diff --git a/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm b/test/tint/builtins/gen/textureLoad/8acf41.wgsl.expected.spvasm index 87554bb606..4ac3e994c3 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 3808a80307..3b57c67a45 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; + + +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. + -******************************************************************** -* 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. * -******************************************************************** diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.hlsl index b3cea20776..4fe76bc1d9 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 ext_tex_plane_1 : register(t2, space1); +cbuffer cbuffer_ext_tex_params : register(b3, space1) { + uint4 ext_tex_params[2]; +}; +Texture2D arg_0 : register(t0, space1); +SamplerState arg_1 : register(s1, space1); +float4 textureSampleExternal(Texture2D plane0, Texture2D 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 21d485b7ec..9b4be131c7 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 -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 plane0, texture2d 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); +} + +void textureSampleLevel_979816(texture2d tint_symbol_1, texture2d 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 tint_symbol_5, texture2d 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 tint_symbol_9 [[texture(0)]], texture2d 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 tint_symbol_13 [[texture(0)]], texture2d 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 tint_symbol_17 [[texture(0)]], texture2d 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; +} -******************************************************************** -* 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. * -******************************************************************** diff --git a/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm b/test/tint/builtins/gen/textureSampleLevel/979816.wgsl.expected.spvasm index 87554bb606..62347b9160 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 0000000000..4a8edf8f8e --- /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) -> vec4 { + return textureLoad(texture, coords); +} + +fn doTextureLoad() { + var res: vec4 = textureLoad2d(arg_0, vec2()); +} + +@stage(vertex) +fn vertex_main() -> @builtin(position) vec4 { + doTextureLoad(); + return vec4(); +} + +@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 0000000000..17d620f537 --- /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 0000000000..e7f6715e89 --- /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 ext_tex_plane_1 : register(t1, space1); +cbuffer cbuffer_ext_tex_params : register(b2, space1) { + uint4 ext_tex_params[2]; +}; +Texture2D arg_0 : register(t0, space1); + +float4 textureLoadExternal(Texture2D plane0, Texture2D 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 tint_symbol, Texture2D 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 0000000000..4e82e4df2a --- /dev/null +++ b/test/tint/builtins/textureLoad/texture_external_param.wgsl.expected.msl @@ -0,0 +1,59 @@ +#include + +using namespace metal; +struct ExternalTextureParams { + /* 0x0000 */ uint numPlanes; + /* 0x0004 */ float vr; + /* 0x0008 */ float ug; + /* 0x000c */ float vg; + /* 0x0010 */ float ub; +}; + +float4 textureLoadExternal(texture2d plane0, texture2d 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 tint_symbol, texture2d 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 tint_symbol_2, texture2d 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 tint_symbol_5, texture2d 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 tint_symbol_8 [[texture(0)]], texture2d 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 tint_symbol_11 [[texture(0)]], texture2d 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 tint_symbol_14 [[texture(0)]], texture2d 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 0000000000..cd8297f6ee --- /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 0000000000..edba45dc64 --- /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) -> vec4 { + return textureLoad(texture, coords); +} + +fn doTextureLoad() { + var res : vec4 = textureLoad2d(arg_0, vec2()); +} + +@stage(vertex) +fn vertex_main() -> @builtin(position) vec4 { + doTextureLoad(); + return vec4(); +} + +@stage(fragment) +fn fragment_main() { + doTextureLoad(); +} + +@stage(compute) @workgroup_size(1) +fn compute_main() { + doTextureLoad(); +}