Convert generators over to use external texture options.
This CL updates the generators to receive the exteral texture options
instead of generating them.
Bug: tint:1855 chromium:1421379
Change-Id: Ib38b902aa441e33d394f947d753075ca6a8fe73b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/123260
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 8d6a3b1..c8d89e4 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -906,8 +906,6 @@
"writer/flatten_bindings.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",
@@ -1627,7 +1625,6 @@
"writer/check_supported_extensions_test.cc",
"writer/flatten_bindings_test.cc",
"writer/float_to_string_test.cc",
- "writer/generate_external_texture_bindings_test.cc",
"writer/text_generator_test.cc",
]
deps = [
@@ -1637,6 +1634,16 @@
]
}
+ # Note, this includes the source files along with the test as otherwise the cmd helpers wouldn't
+ # be included in the binary.
+ tint_unittests_source_set("tint_unittests_cmd_src") {
+ sources = [
+ "cmd/generate_external_texture_bindings.cc",
+ "cmd/generate_external_texture_bindings.h",
+ "cmd/generate_external_texture_bindings_test.cc",
+ ]
+ }
+
tint_unittests_source_set("tint_unittests_spv_reader_src") {
sources = [
"reader/spirv/enum_converter_test.cc",
@@ -2033,6 +2040,7 @@
":tint_unittests_ast_src",
":tint_unittests_base_src",
":tint_unittests_builtins_src",
+ ":tint_unittests_cmd_src",
":tint_unittests_constant_src",
":tint_unittests_core_src",
":tint_unittests_demangler_src",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 0f442ff..0c2da7a 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -550,8 +550,6 @@
writer/flatten_bindings.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
@@ -1022,10 +1020,17 @@
writer/check_supported_extensions_test.cc
writer/flatten_bindings_test.cc
writer/float_to_string_test.cc
- writer/generate_external_texture_bindings_test.cc
writer/text_generator_test.cc
)
+ # Noet, the source files are included here otherwise the cmd sources would not be included in the
+ # test binary.
+ list(APPEND TINT_TEST_SRCS
+ cmd/generate_external_texture_bindings.cc
+ cmd/generate_external_texture_bindings.h
+ cmd/generate_external_texture_bindings_test.cc
+ )
+
# Uniformity analysis tests depend on WGSL reader
if(${TINT_BUILD_WGSL_READER})
list(APPEND TINT_TEST_SRCS
diff --git a/src/tint/cmd/BUILD.gn b/src/tint/cmd/BUILD.gn
index b4ff6d5..946bd36 100644
--- a/src/tint/cmd/BUILD.gn
+++ b/src/tint/cmd/BUILD.gn
@@ -17,6 +17,8 @@
source_set("tint_cmd_helper") {
sources = [
+ "generate_external_texture_bindings.cc",
+ "generate_external_texture_bindings.h",
"helper.cc",
"helper.h",
]
diff --git a/src/tint/cmd/CMakeLists.txt b/src/tint/cmd/CMakeLists.txt
index 6da9701..142595b 100644
--- a/src/tint/cmd/CMakeLists.txt
+++ b/src/tint/cmd/CMakeLists.txt
@@ -15,6 +15,8 @@
## Tint executable
add_executable(tint "")
target_sources(tint PRIVATE
+ "generate_external_texture_bindings.cc"
+ "generate_external_texture_bindings.h"
"helper.cc"
"helper.h"
"main.cc"
diff --git a/src/tint/writer/generate_external_texture_bindings.cc b/src/tint/cmd/generate_external_texture_bindings.cc
similarity index 82%
rename from src/tint/writer/generate_external_texture_bindings.cc
rename to src/tint/cmd/generate_external_texture_bindings.cc
index 6e5daf0..1b6ca74 100644
--- a/src/tint/writer/generate_external_texture_bindings.cc
+++ b/src/tint/cmd/generate_external_texture_bindings.cc
@@ -12,7 +12,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/writer/generate_external_texture_bindings.h"
+#include "src/tint/cmd/generate_external_texture_bindings.h"
#include <algorithm>
#include <unordered_map>
@@ -20,12 +20,13 @@
#include "src/tint/ast/module.h"
#include "src/tint/program.h"
+#include "src/tint/sem/binding_point.h"
#include "src/tint/sem/variable.h"
#include "src/tint/type/external_texture.h"
-namespace tint::writer {
+namespace tint::cmd {
-transform::MultiplanarExternalTexture::BindingsMap GenerateExternalTextureBindings(
+writer::ExternalTextureOptions::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.
@@ -45,15 +46,17 @@
}
}
- transform::MultiplanarExternalTexture::BindingsMap new_bindings_map;
+ writer::ExternalTextureOptions::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::MultiplanarExternalTexture::BindingPoints{{g, next_num++}, {g, next_num++}};
+ writer::ExternalTextureOptions::BindingPoints{{g, next_num++}, {g, next_num++}};
+
new_bindings_map[bp] = new_bps;
}
+
return new_bindings_map;
}
-} // namespace tint::writer
+} // namespace tint::cmd
diff --git a/src/tint/cmd/generate_external_texture_bindings.h b/src/tint/cmd/generate_external_texture_bindings.h
new file mode 100644
index 0000000..c6a842d
--- /dev/null
+++ b/src/tint/cmd/generate_external_texture_bindings.h
@@ -0,0 +1,26 @@
+// 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_CMD_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
+#define SRC_TINT_CMD_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
+
+#include "tint/tint.h"
+
+namespace tint::cmd {
+
+writer::ExternalTextureOptions::BindingsMap GenerateExternalTextureBindings(const Program* program);
+
+} // namespace tint::cmd
+
+#endif // SRC_TINT_CMD_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_
diff --git a/src/tint/writer/generate_external_texture_bindings_test.cc b/src/tint/cmd/generate_external_texture_bindings_test.cc
similarity index 97%
rename from src/tint/writer/generate_external_texture_bindings_test.cc
rename to src/tint/cmd/generate_external_texture_bindings_test.cc
index b862554..2c19675 100644
--- a/src/tint/writer/generate_external_texture_bindings_test.cc
+++ b/src/tint/cmd/generate_external_texture_bindings_test.cc
@@ -15,11 +15,11 @@
#include <utility>
#include "gtest/gtest.h"
+#include "src/tint/cmd/generate_external_texture_bindings.h"
#include "src/tint/program_builder.h"
#include "src/tint/writer/binding_point.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
-namespace tint::writer {
+namespace tint::cmd {
namespace {
using namespace tint::number_suffixes; // NOLINT
@@ -126,4 +126,4 @@
}
} // namespace
-} // namespace tint::writer
+} // namespace tint::cmd
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index f4ad19f..2daf75f 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -39,6 +39,7 @@
#endif // TINT_BUILD_SPV_READER
#include "src/tint/ast/module.h"
+#include "src/tint/cmd/generate_external_texture_bindings.h"
#include "src/tint/cmd/helper.h"
#include "src/tint/utils/io/command.h"
#include "src/tint/utils/string.h"
@@ -549,7 +550,8 @@
tint::writer::spirv::Options gen_options;
gen_options.disable_robustness = !options.enable_robustness;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
- gen_options.generate_external_texture_bindings = true;
+ gen_options.external_texture_options.bindings_map =
+ tint::cmd::GenerateExternalTextureBindings(program);
auto result = tint::writer::spirv::Generate(program, gen_options);
if (!result.success) {
tint::cmd::PrintWGSL(std::cerr, *program);
@@ -656,7 +658,8 @@
tint::writer::msl::Options gen_options;
gen_options.disable_robustness = !options.enable_robustness;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
- gen_options.generate_external_texture_bindings = true;
+ gen_options.external_texture_options.bindings_map =
+ tint::cmd::GenerateExternalTextureBindings(input_program);
auto result = tint::writer::msl::Generate(input_program, gen_options);
if (!result.success) {
tint::cmd::PrintWGSL(std::cerr, *program);
@@ -717,7 +720,8 @@
tint::writer::hlsl::Options gen_options;
gen_options.disable_robustness = !options.enable_robustness;
gen_options.disable_workgroup_init = options.disable_workgroup_init;
- gen_options.generate_external_texture_bindings = true;
+ gen_options.external_texture_options.bindings_map =
+ tint::cmd::GenerateExternalTextureBindings(program);
gen_options.root_constant_binding_point = options.hlsl_root_constant_binding_point;
auto result = tint::writer::hlsl::Generate(program, gen_options);
if (!result.success) {
@@ -856,7 +860,8 @@
auto generate = [&](const tint::Program* prg, const std::string entry_point_name) -> bool {
tint::writer::glsl::Options gen_options;
gen_options.disable_robustness = !options.enable_robustness;
- gen_options.generate_external_texture_bindings = true;
+ gen_options.external_texture_options.bindings_map =
+ tint::cmd::GenerateExternalTextureBindings(prg);
auto result = tint::writer::glsl::Generate(prg, gen_options, entry_point_name);
if (!result.success) {
tint::cmd::PrintWGSL(std::cerr, *prg);
@@ -1003,55 +1008,6 @@
m.Add<tint::transform::SubstituteOverride>();
return true;
}},
- {"multiplaner_external_texture",
- [](tint::inspector::Inspector& inspector, tint::transform::Manager& m,
- tint::transform::DataMap& i) {
- using MET = tint::transform::MultiplanarExternalTexture;
-
- // Generate the MultiplanarExternalTexture::NewBindingPoints by finding two free
- // binding points. We may wish to expose these binding points via a command line flag
- // in the future.
-
- // Set of all the group-0 bindings in use.
- std::unordered_set<uint32_t> group0_bindings_in_use;
- auto allocate_binding = [&] {
- for (uint32_t idx = 0;; idx++) {
- auto binding = tint::writer::BindingPoint{0u, idx};
- if (group0_bindings_in_use.emplace(idx).second) {
- return binding;
- }
- }
- };
- // Populate group0_bindings_in_use with the existing bindings across all entry points.
- for (auto ep : inspector.GetEntryPoints()) {
- for (auto binding : inspector.GetResourceBindings(ep.name)) {
- if (binding.bind_group == 0) {
- group0_bindings_in_use.emplace(binding.binding);
- }
- }
- }
- // Allocate new binding points for the external texture's planes and parameters.
- MET::BindingsMap met_bindings;
- for (auto ep : inspector.GetEntryPoints()) {
- for (auto ext_tex : inspector.GetExternalTextureResourceBindings(ep.name)) {
- auto binding = tint::writer::BindingPoint{
- ext_tex.bind_group,
- ext_tex.binding,
- };
- if (met_bindings.count(binding)) {
- continue;
- }
- met_bindings.emplace(binding, MET::BindingPoints{
- /* plane_1 */ allocate_binding(),
- /* params */ allocate_binding(),
- });
- }
- }
-
- i.Add<MET::NewBindingPoints>(std::move(met_bindings));
- m.Add<MET>();
- return true;
- }},
};
auto transform_names = [&] {
tint::utils::StringStream names;
diff --git a/src/tint/writer/generate_external_texture_bindings.h b/src/tint/writer/generate_external_texture_bindings.h
deleted file mode 100644
index 8d0aad9..0000000
--- a/src/tint/writer/generate_external_texture_bindings.h
+++ /dev/null
@@ -1,27 +0,0 @@
-// 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/glsl/generator.h b/src/tint/writer/glsl/generator.h
index 5c2de66..ef7b65b 100644
--- a/src/tint/writer/glsl/generator.h
+++ b/src/tint/writer/glsl/generator.h
@@ -74,9 +74,6 @@
/// 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 in the binding mappings for external textures
ExternalTextureOptions external_texture_options = {};
@@ -87,7 +84,6 @@
TINT_REFLECT(disable_robustness,
allow_collisions,
disable_workgroup_init,
- generate_external_texture_bindings,
external_texture_options,
version);
};
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 29c28ff..cdf4a15 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -53,6 +53,7 @@
#include "src/tint/transform/disable_uniformity_analysis.h"
#include "src/tint/transform/expand_compound_assignment.h"
#include "src/tint/transform/manager.h"
+#include "src/tint/transform/multiplanar_external_texture.h"
#include "src/tint/transform/pad_structs.h"
#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_initializers_to_let.h"
@@ -81,7 +82,6 @@
#include "src/tint/utils/string_stream.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"
using namespace tint::number_suffixes; // NOLINT
@@ -180,10 +180,10 @@
manager.Add<transform::Robustness>();
}
- if (options.generate_external_texture_bindings) {
+ if (!options.external_texture_options.bindings_map.empty()) {
// Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
- auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
- data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ options.external_texture_options.bindings_map);
manager.Add<transform::MultiplanarExternalTexture>();
}
diff --git a/src/tint/writer/hlsl/generator.h b/src/tint/writer/hlsl/generator.h
index 1c072d8..2cdfe6f 100644
--- a/src/tint/writer/hlsl/generator.h
+++ b/src/tint/writer/hlsl/generator.h
@@ -59,9 +59,6 @@
/// 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 in the binding mappings for external textures
ExternalTextureOptions external_texture_options = {};
@@ -80,7 +77,6 @@
TINT_REFLECT(disable_robustness,
root_constant_binding_point,
disable_workgroup_init,
- generate_external_texture_bindings,
external_texture_options,
array_length_from_uniform);
};
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 84543ab..1cfbf64 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -53,6 +53,7 @@
#include "src/tint/transform/expand_compound_assignment.h"
#include "src/tint/transform/localize_struct_array_assignment.h"
#include "src/tint/transform/manager.h"
+#include "src/tint/transform/multiplanar_external_texture.h"
#include "src/tint/transform/num_workgroups_from_uniform.h"
#include "src/tint/transform/promote_initializers_to_let.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
@@ -81,7 +82,6 @@
#include "src/tint/writer/append_vector.h"
#include "src/tint/writer/check_supported_extensions.h"
#include "src/tint/writer/float_to_string.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
using namespace tint::number_suffixes; // NOLINT
@@ -188,10 +188,10 @@
manager.Add<transform::Robustness>();
}
- if (options.generate_external_texture_bindings) {
+ if (!options.external_texture_options.bindings_map.empty()) {
// Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
- auto new_bindings_map = GenerateExternalTextureBindings(in);
- data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ options.external_texture_options.bindings_map);
manager.Add<transform::MultiplanarExternalTexture>();
}
diff --git a/src/tint/writer/msl/generator.h b/src/tint/writer/msl/generator.h
index e58ef9e..b4e11d9 100644
--- a/src/tint/writer/msl/generator.h
+++ b/src/tint/writer/msl/generator.h
@@ -63,9 +63,6 @@
/// 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 in the binding mappings for external textures
ExternalTextureOptions external_texture_options = {};
@@ -79,7 +76,6 @@
fixed_sample_mask,
emit_vertex_point_size,
disable_workgroup_init,
- generate_external_texture_bindings,
external_texture_options,
array_length_from_uniform);
};
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 397c961..268736f 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -49,6 +49,7 @@
#include "src/tint/transform/expand_compound_assignment.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/module_scope_var_to_entry_point_param.h"
+#include "src/tint/transform/multiplanar_external_texture.h"
#include "src/tint/transform/packed_vec3.h"
#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_initializers_to_let.h"
@@ -83,7 +84,6 @@
#include "src/tint/utils/string_stream.h"
#include "src/tint/writer/check_supported_extensions.h"
#include "src/tint/writer/float_to_string.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint::writer::msl {
namespace {
@@ -232,10 +232,10 @@
manager.Add<transform::BuiltinPolyfill>();
}
- if (options.generate_external_texture_bindings) {
+ if (!options.external_texture_options.bindings_map.empty()) {
// Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
- auto new_bindings_map = GenerateExternalTextureBindings(in);
- data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ options.external_texture_options.bindings_map);
manager.Add<transform::MultiplanarExternalTexture>();
}
diff --git a/src/tint/writer/spirv/generator.h b/src/tint/writer/spirv/generator.h
index 11660fe..5ad9057 100644
--- a/src/tint/writer/spirv/generator.h
+++ b/src/tint/writer/spirv/generator.h
@@ -46,9 +46,6 @@
/// 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 in the binding mappings for external textures
ExternalTextureOptions external_texture_options = {};
@@ -60,7 +57,6 @@
TINT_REFLECT(disable_robustness,
emit_vertex_point_size,
disable_workgroup_init,
- generate_external_texture_bindings,
external_texture_options,
use_zero_initialize_workgroup_memory_extension);
};
diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
index 3c74920..f3bbaf3 100644
--- a/src/tint/writer/spirv/generator_impl.cc
+++ b/src/tint/writer/spirv/generator_impl.cc
@@ -28,6 +28,7 @@
#include "src/tint/transform/for_loop_to_loop.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/merge_return.h"
+#include "src/tint/transform/multiplanar_external_texture.h"
#include "src/tint/transform/preserve_padding.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -41,7 +42,6 @@
#include "src/tint/transform/vectorize_scalar_matrix_initializers.h"
#include "src/tint/transform/while_to_loop.h"
#include "src/tint/transform/zero_init_workgroup_memory.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint::writer::spirv {
@@ -73,10 +73,10 @@
manager.Add<transform::Robustness>();
}
- if (options.generate_external_texture_bindings) {
+ if (!options.external_texture_options.bindings_map.empty()) {
// Note: it is more efficient for MultiplanarExternalTexture to come after Robustness
- auto new_bindings_map = GenerateExternalTextureBindings(in);
- data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ options.external_texture_options.bindings_map);
manager.Add<transform::MultiplanarExternalTexture>();
}
diff --git a/test/tint/bug/tint/1739.wgsl b/test/tint/bug/tint/1739.wgsl
index d513d64..ef45dd5 100644
--- a/test/tint/bug/tint/1739.wgsl
+++ b/test/tint/bug/tint/1739.wgsl
@@ -1,4 +1,4 @@
-// flags: --transform robustness,multiplaner_external_texture
+// flags: --transform robustness
@group(0) @binding(0) var t : texture_external;
@group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>;
diff --git a/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl
index e71fad6..27ce96e 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl
@@ -31,33 +31,18 @@
float3 gammaCorrection(float3 v, GammaTransferParams params) {
const bool3 cond = (abs(v) < float3((params.D).xxx));
- const float3 t_1 = (float3(sign(v)) * ((params.C * abs(v)) + params.F));
+ const float3 t = (float3(sign(v)) * ((params.C * abs(v)) + params.F));
const float3 f = (float3(sign(v)) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
- return (cond ? t_1 : f);
+ return (cond ? t : f);
}
float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, int2 coord, ExternalTextureParams params) {
const int2 coord1 = (coord >> (1u).xx);
float3 color = float3(0.0f, 0.0f, 0.0f);
if ((params.numPlanes == 1u)) {
- uint3 tint_tmp;
- plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
- const uint level_idx = min(0u, (tint_tmp.z - 1u));
- uint3 tint_tmp_1;
- plane0.GetDimensions(level_idx, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
- color = plane0.Load(int3(tint_clamp(coord, (0).xx, int2((tint_tmp_1.xy - (1u).xx))), int(level_idx))).rgb;
+ color = plane0.Load(int3(coord, 0)).rgb;
} else {
- uint3 tint_tmp_2;
- plane0.GetDimensions(0, tint_tmp_2.x, tint_tmp_2.y, tint_tmp_2.z);
- const uint level_idx_1 = min(0u, (tint_tmp_2.z - 1u));
- uint3 tint_tmp_3;
- plane1.GetDimensions(0, tint_tmp_3.x, tint_tmp_3.y, tint_tmp_3.z);
- const uint level_idx_2 = min(0u, (tint_tmp_3.z - 1u));
- uint3 tint_tmp_4;
- plane0.GetDimensions(level_idx_1, tint_tmp_4.x, tint_tmp_4.y, tint_tmp_4.z);
- uint3 tint_tmp_5;
- plane1.GetDimensions(level_idx_2, tint_tmp_5.x, tint_tmp_5.y, tint_tmp_5.z);
- color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(tint_clamp(coord, (0).xx, int2((tint_tmp_4.xy - (1u).xx))), int(level_idx_1))).r, plane1.Load(int3(tint_clamp(coord1, (0).xx, int2((tint_tmp_5.xy - (1u).xx))), int(level_idx_2))).rg, 1.0f));
+ color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord1, 0)).rg, 1.0f));
}
if ((params.doYuvToRgbConversionOnly == 0u)) {
color = gammaCorrection(color, params.gammaDecodeParams);
@@ -113,13 +98,17 @@
[numthreads(1, 1, 1)]
void main() {
- float4 red = textureLoadExternal(t, ext_tex_plane_1, (10).xx, ext_tex_params_load(0u));
- uint2 tint_tmp_6;
- outImage.GetDimensions(tint_tmp_6.x, tint_tmp_6.y);
- outImage[tint_clamp((0).xx, (0).xx, int2((tint_tmp_6 - (1u).xx)))] = red;
- float4 green = textureLoadExternal(t, ext_tex_plane_1, int2(70, 118), ext_tex_params_load(0u));
- uint2 tint_tmp_7;
- outImage.GetDimensions(tint_tmp_7.x, tint_tmp_7.y);
- outImage[tint_clamp(int2(1, 0), (0).xx, int2((tint_tmp_7 - (1u).xx)))] = green;
+ uint2 tint_tmp;
+ t.GetDimensions(tint_tmp.x, tint_tmp.y);
+ float4 red = textureLoadExternal(t, ext_tex_plane_1, tint_clamp((10).xx, (0).xx, int2((tint_tmp - (1u).xx))), ext_tex_params_load(0u));
+ uint2 tint_tmp_1;
+ outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
+ outImage[tint_clamp((0).xx, (0).xx, int2((tint_tmp_1 - (1u).xx)))] = red;
+ uint2 tint_tmp_2;
+ t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
+ float4 green = textureLoadExternal(t, ext_tex_plane_1, tint_clamp(int2(70, 118), (0).xx, int2((tint_tmp_2 - (1u).xx))), ext_tex_params_load(0u));
+ uint2 tint_tmp_3;
+ outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y);
+ outImage[tint_clamp(int2(1, 0), (0).xx, int2((tint_tmp_3 - (1u).xx)))] = green;
return;
}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl
index e71fad6..27ce96e 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl
@@ -31,33 +31,18 @@
float3 gammaCorrection(float3 v, GammaTransferParams params) {
const bool3 cond = (abs(v) < float3((params.D).xxx));
- const float3 t_1 = (float3(sign(v)) * ((params.C * abs(v)) + params.F));
+ const float3 t = (float3(sign(v)) * ((params.C * abs(v)) + params.F));
const float3 f = (float3(sign(v)) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E));
- return (cond ? t_1 : f);
+ return (cond ? t : f);
}
float4 textureLoadExternal(Texture2D<float4> plane0, Texture2D<float4> plane1, int2 coord, ExternalTextureParams params) {
const int2 coord1 = (coord >> (1u).xx);
float3 color = float3(0.0f, 0.0f, 0.0f);
if ((params.numPlanes == 1u)) {
- uint3 tint_tmp;
- plane0.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
- const uint level_idx = min(0u, (tint_tmp.z - 1u));
- uint3 tint_tmp_1;
- plane0.GetDimensions(level_idx, tint_tmp_1.x, tint_tmp_1.y, tint_tmp_1.z);
- color = plane0.Load(int3(tint_clamp(coord, (0).xx, int2((tint_tmp_1.xy - (1u).xx))), int(level_idx))).rgb;
+ color = plane0.Load(int3(coord, 0)).rgb;
} else {
- uint3 tint_tmp_2;
- plane0.GetDimensions(0, tint_tmp_2.x, tint_tmp_2.y, tint_tmp_2.z);
- const uint level_idx_1 = min(0u, (tint_tmp_2.z - 1u));
- uint3 tint_tmp_3;
- plane1.GetDimensions(0, tint_tmp_3.x, tint_tmp_3.y, tint_tmp_3.z);
- const uint level_idx_2 = min(0u, (tint_tmp_3.z - 1u));
- uint3 tint_tmp_4;
- plane0.GetDimensions(level_idx_1, tint_tmp_4.x, tint_tmp_4.y, tint_tmp_4.z);
- uint3 tint_tmp_5;
- plane1.GetDimensions(level_idx_2, tint_tmp_5.x, tint_tmp_5.y, tint_tmp_5.z);
- color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(tint_clamp(coord, (0).xx, int2((tint_tmp_4.xy - (1u).xx))), int(level_idx_1))).r, plane1.Load(int3(tint_clamp(coord1, (0).xx, int2((tint_tmp_5.xy - (1u).xx))), int(level_idx_2))).rg, 1.0f));
+ color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord1, 0)).rg, 1.0f));
}
if ((params.doYuvToRgbConversionOnly == 0u)) {
color = gammaCorrection(color, params.gammaDecodeParams);
@@ -113,13 +98,17 @@
[numthreads(1, 1, 1)]
void main() {
- float4 red = textureLoadExternal(t, ext_tex_plane_1, (10).xx, ext_tex_params_load(0u));
- uint2 tint_tmp_6;
- outImage.GetDimensions(tint_tmp_6.x, tint_tmp_6.y);
- outImage[tint_clamp((0).xx, (0).xx, int2((tint_tmp_6 - (1u).xx)))] = red;
- float4 green = textureLoadExternal(t, ext_tex_plane_1, int2(70, 118), ext_tex_params_load(0u));
- uint2 tint_tmp_7;
- outImage.GetDimensions(tint_tmp_7.x, tint_tmp_7.y);
- outImage[tint_clamp(int2(1, 0), (0).xx, int2((tint_tmp_7 - (1u).xx)))] = green;
+ uint2 tint_tmp;
+ t.GetDimensions(tint_tmp.x, tint_tmp.y);
+ float4 red = textureLoadExternal(t, ext_tex_plane_1, tint_clamp((10).xx, (0).xx, int2((tint_tmp - (1u).xx))), ext_tex_params_load(0u));
+ uint2 tint_tmp_1;
+ outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
+ outImage[tint_clamp((0).xx, (0).xx, int2((tint_tmp_1 - (1u).xx)))] = red;
+ uint2 tint_tmp_2;
+ t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
+ float4 green = textureLoadExternal(t, ext_tex_plane_1, tint_clamp(int2(70, 118), (0).xx, int2((tint_tmp_2 - (1u).xx))), ext_tex_params_load(0u));
+ uint2 tint_tmp_3;
+ outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y);
+ outImage[tint_clamp(int2(1, 0), (0).xx, int2((tint_tmp_3 - (1u).xx)))] = green;
return;
}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.glsl b/test/tint/bug/tint/1739.wgsl.expected.glsl
index 3bbb840..c32e19b 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.glsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.glsl
@@ -1,7 +1,10 @@
-SKIP: FAILED
-
#version 310 es
+vec3 tint_select(vec3 param_0, vec3 param_1, bvec3 param_2) {
+ return vec3(param_2[0] ? param_1[0] : param_0[0], param_2[1] ? param_1[1] : param_0[1], param_2[2] ? param_1[2] : param_0[2]);
+}
+
+
struct GammaTransferParams {
float G;
float A;
@@ -50,18 +53,18 @@
layout(rgba8) uniform highp writeonly image2D outImage;
vec3 gammaCorrection(vec3 v, GammaTransferParams params) {
bvec3 cond = lessThan(abs(v), vec3(params.D));
- vec3 t_1 = (sign(v) * ((params.C * abs(v)) + params.F));
+ vec3 t = (sign(v) * ((params.C * abs(v)) + params.F));
vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E));
- return mix(f, t_1, cond);
+ return tint_select(f, t, cond);
}
vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) {
ivec2 coord1 = (coord >> uvec2(1u));
vec3 color = vec3(0.0f, 0.0f, 0.0f);
if ((params.numPlanes == 1u)) {
- color = texelFetch(plane0_1, clamp(coord, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane0_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))).rgb;
+ color = texelFetch(plane0_1, coord, 0).rgb;
} else {
- color = (vec4(texelFetch(plane0_1, clamp(coord, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane0_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane0_1))) - 1u)))).r, texelFetch(plane1_1, clamp(coord1, ivec2(0), ivec2((uvec2(uvec2(textureSize(plane1_1, clamp(0, 0, int((uint(uint(textureQueryLevels(plane1_1))) - 1u)))))) - uvec2(1u)))), clamp(0, 0, int((uint(uint(textureQueryLevels(plane1_1))) - 1u)))).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+ color = (vec4(texelFetch(plane0_1, coord, 0).r, texelFetch(plane1_1, coord1, 0).rg, 1.0f) * params.yuvToRgbConversionMatrix);
}
if ((params.doYuvToRgbConversionOnly == 0u)) {
color = gammaCorrection(color, params.gammaDecodeParams);
@@ -71,17 +74,17 @@
return vec4(color, 1.0f);
}
-uniform highp sampler2D t_2;
+uniform highp sampler2D t_1;
uniform highp sampler2D ext_tex_plane_1_1;
ExternalTextureParams conv_ExternalTextureParams(ExternalTextureParams_std140 val) {
return ExternalTextureParams(val.numPlanes, val.doYuvToRgbConversionOnly, val.pad, val.pad_1, val.yuvToRgbConversionMatrix, val.gammaDecodeParams, val.gammaEncodeParams, val.gamutConversionMatrix, mat3x2(val.coordTransformationMatrix_0, val.coordTransformationMatrix_1, val.coordTransformationMatrix_2), val.pad_2, val.pad_3);
}
void tint_symbol() {
- vec4 red = textureLoadExternal(t_2, ext_tex_plane_1_1, ivec2(10), conv_ExternalTextureParams(ext_tex_params.inner));
- imageStore(outImage, clamp(ivec2(0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), red);
- vec4 green = textureLoadExternal(t_2, ext_tex_plane_1_1, ivec2(70, 118), conv_ExternalTextureParams(ext_tex_params.inner));
- imageStore(outImage, clamp(ivec2(1, 0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), green);
+ vec4 red = textureLoadExternal(t_1, ext_tex_plane_1_1, clamp(ivec2(10), ivec2(0), ivec2((uvec2(textureSize(t_1, 0)) - uvec2(1u)))), conv_ExternalTextureParams(ext_tex_params.inner));
+ imageStore(outImage, clamp(ivec2(0), ivec2(0), ivec2((uvec2(imageSize(outImage)) - uvec2(1u)))), red);
+ vec4 green = textureLoadExternal(t_1, ext_tex_plane_1_1, clamp(ivec2(70, 118), ivec2(0), ivec2((uvec2(textureSize(t_1, 0)) - uvec2(1u)))), conv_ExternalTextureParams(ext_tex_params.inner));
+ imageStore(outImage, clamp(ivec2(1, 0), ivec2(0), ivec2((uvec2(imageSize(outImage)) - uvec2(1u)))), green);
return;
}
@@ -90,10 +93,3 @@
tint_symbol();
return;
}
-Error parsing GLSL shader:
-ERROR: 0:60: 'textureQueryLevels' : no matching overloaded function found
-ERROR: 0:60: '' : compilation terminated
-ERROR: 2 compilation errors. No code generated.
-
-
-
diff --git a/test/tint/bug/tint/1739.wgsl.expected.msl b/test/tint/bug/tint/1739.wgsl.expected.msl
index 9bcbc79..34134b3 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.msl
+++ b/test/tint/bug/tint/1739.wgsl.expected.msl
@@ -78,21 +78,18 @@
float3 gammaCorrection(float3 v, GammaTransferParams params) {
bool3 const cond = (fabs(v) < float3(params.D));
- float3 const t_1 = (sign(v) * ((params.C * fabs(v)) + params.F));
+ float3 const t = (sign(v) * ((params.C * fabs(v)) + params.F));
float3 const f = (sign(v) * (pow(((params.A * fabs(v)) + params.B), float3(params.G)) + params.E));
- return select(f, t_1, cond);
+ return select(f, t, cond);
}
float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<float, access::sample> plane1, int2 coord, ExternalTextureParams params) {
int2 const coord1 = (coord >> uint2(1u));
float3 color = 0.0f;
if ((params.numPlanes == 1u)) {
- uint const level_idx = min(0u, (plane0.get_num_mip_levels() - 1u));
- color = plane0.read(uint2(tint_clamp(coord, int2(0), int2((uint2(plane0.get_width(level_idx), plane0.get_height(level_idx)) - uint2(1u))))), level_idx).rgb;
+ color = plane0.read(uint2(coord), 0).rgb;
} else {
- uint const level_idx_1 = min(0u, (plane0.get_num_mip_levels() - 1u));
- uint const level_idx_2 = min(0u, (plane1.get_num_mip_levels() - 1u));
- color = (float4(plane0.read(uint2(tint_clamp(coord, int2(0), int2((uint2(plane0.get_width(level_idx_1), plane0.get_height(level_idx_1)) - uint2(1u))))), level_idx_1)[0], plane1.read(uint2(tint_clamp(coord1, int2(0), int2((uint2(plane1.get_width(level_idx_2), plane1.get_height(level_idx_2)) - uint2(1u))))), level_idx_2).rg, 1.0f) * params.yuvToRgbConversionMatrix);
+ color = (float4(plane0.read(uint2(coord), 0)[0], plane1.read(uint2(coord1), 0).rg, 1.0f) * params.yuvToRgbConversionMatrix);
}
if ((params.doYuvToRgbConversionOnly == 0u)) {
color = gammaCorrection(color, params.gammaDecodeParams);
@@ -102,10 +99,10 @@
return float4(color, 1.0f);
}
-kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_1 [[texture(0)]], texture2d<float, access::sample> tint_symbol_2 [[texture(1)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_3 [[buffer(0)]], texture2d<float, access::write> tint_symbol_4 [[texture(2)]]) {
- float4 red = textureLoadExternal(tint_symbol_1, tint_symbol_2, int2(10), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_1 [[texture(1)]], texture2d<float, access::sample> tint_symbol_2 [[texture(2)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_3 [[buffer(3)]], texture2d<float, access::write> tint_symbol_4 [[texture(0)]]) {
+ float4 red = textureLoadExternal(tint_symbol_1, tint_symbol_2, tint_clamp(int2(10), int2(0), int2((uint2(tint_symbol_1.get_width(), tint_symbol_1.get_height()) - uint2(1u)))), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
tint_symbol_4.write(red, uint2(tint_clamp(int2(0), int2(0), int2((uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height()) - uint2(1u))))));
- float4 green = textureLoadExternal(tint_symbol_1, tint_symbol_2, int2(70, 118), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
+ float4 green = textureLoadExternal(tint_symbol_1, tint_symbol_2, tint_clamp(int2(70, 118), int2(0), int2((uint2(tint_symbol_1.get_width(), tint_symbol_1.get_height()) - uint2(1u)))), tint_unpack_vec3_in_composite_1(*(tint_symbol_3)));
tint_symbol_4.write(green, uint2(tint_clamp(int2(1, 0), int2(0), int2((uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height()) - uint2(1u))))));
return;
}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.spvasm b/test/tint/bug/tint/1739.wgsl.expected.spvasm
index b31df45..01e7c86 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/1739.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 203
+; Bound: 193
; Schema: 0
OpCapability Shader
OpCapability ImageQuery
@@ -139,24 +139,25 @@
%v2uint = OpTypeVector %uint 2
%uint_1 = OpConstant %uint 1
%81 = OpConstantComposite %v2uint %uint_1 %uint_1
- %90 = OpConstantNull %uint
- %95 = OpConstantNull %v2int
+ %90 = OpConstantNull %int
%float_1 = OpConstant %float 1
- %142 = OpTypeFunction %ExternalTextureParams %ExternalTextureParams_std140
+ %103 = OpConstantNull %uint
+ %121 = OpTypeFunction %ExternalTextureParams %ExternalTextureParams_std140
%void = OpTypeVoid
- %157 = OpTypeFunction %void
+ %136 = OpTypeFunction %void
%int_10 = OpConstant %int 10
- %165 = OpConstantComposite %v2int %int_10 %int_10
+ %145 = OpConstantComposite %v2int %int_10 %int_10
+ %146 = OpConstantNull %v2int
+ %int_0 = OpConstant %int 0
%uint_0 = OpConstant %uint 0
%_ptr_Uniform_ExternalTextureParams_std140 = OpTypePointer Uniform %ExternalTextureParams_std140
%_ptr_Function_v4float = OpTypePointer Function %v4float
- %173 = OpConstantNull %v4float
+ %159 = OpConstantNull %v4float
%int_70 = OpConstant %int 70
%int_118 = OpConstant %int 118
- %187 = OpConstantComposite %v2int %int_70 %int_118
+ %174 = OpConstantComposite %v2int %int_70 %int_118
%int_1 = OpConstant %int 1
- %196 = OpConstantNull %int
- %197 = OpConstantComposite %v2int %int_1 %196
+ %187 = OpConstantComposite %v2int %int_1 %90
%tint_clamp = OpFunction %v2int None %20
%e = OpFunctionParameter %v2int
%low = OpFunctionParameter %v2int
@@ -215,119 +216,108 @@
OpSelectionMerge %86 None
OpBranchConditional %85 %87 %88
%87 = OpLabel
- %91 = OpImageQueryLevels %uint %plane0
- %92 = OpISub %uint %91 %uint_1
- %89 = OpExtInst %uint %29 UMin %90 %92
- %97 = OpImageQuerySizeLod %v2uint %plane0 %89
- %98 = OpISub %v2uint %97 %81
- %96 = OpBitcast %v2int %98
- %94 = OpFunctionCall %v2int %tint_clamp %coord %95 %96
- %93 = OpImageFetch %v4float %plane0 %94 Lod %89
- %99 = OpVectorShuffle %v3float %93 %93 0 1 2
- OpStore %color %99
+ %89 = OpImageFetch %v4float %plane0 %coord Lod %90
+ %91 = OpVectorShuffle %v3float %89 %89 0 1 2
+ OpStore %color %91
OpBranch %86
%88 = OpLabel
- %101 = OpImageQueryLevels %uint %plane0
- %102 = OpISub %uint %101 %uint_1
- %100 = OpExtInst %uint %29 UMin %90 %102
- %104 = OpImageQueryLevels %uint %plane1
- %105 = OpISub %uint %104 %uint_1
- %103 = OpExtInst %uint %29 UMin %90 %105
- %109 = OpImageQuerySizeLod %v2uint %plane0 %100
- %110 = OpISub %v2uint %109 %81
- %108 = OpBitcast %v2int %110
- %107 = OpFunctionCall %v2int %tint_clamp %coord %95 %108
- %106 = OpImageFetch %v4float %plane0 %107 Lod %100
- %111 = OpCompositeExtract %float %106 0
- %115 = OpImageQuerySizeLod %v2uint %plane1 %103
- %116 = OpISub %v2uint %115 %81
- %114 = OpBitcast %v2int %116
- %113 = OpFunctionCall %v2int %tint_clamp %82 %95 %114
- %112 = OpImageFetch %v4float %plane1 %113 Lod %103
- %117 = OpVectorShuffle %v2float %112 %112 0 1
- %118 = OpCompositeExtract %float %117 0
- %119 = OpCompositeExtract %float %117 1
- %121 = OpCompositeConstruct %v4float %111 %118 %119 %float_1
- %122 = OpCompositeExtract %mat3v4float %params_0 2
- %123 = OpVectorTimesMatrix %v3float %121 %122
- OpStore %color %123
+ %92 = OpImageFetch %v4float %plane0 %coord Lod %90
+ %93 = OpCompositeExtract %float %92 0
+ %94 = OpImageFetch %v4float %plane1 %82 Lod %90
+ %95 = OpVectorShuffle %v2float %94 %94 0 1
+ %96 = OpCompositeExtract %float %95 0
+ %97 = OpCompositeExtract %float %95 1
+ %99 = OpCompositeConstruct %v4float %93 %96 %97 %float_1
+ %100 = OpCompositeExtract %mat3v4float %params_0 2
+ %101 = OpVectorTimesMatrix %v3float %99 %100
+ OpStore %color %101
OpBranch %86
%86 = OpLabel
- %124 = OpCompositeExtract %uint %params_0 1
- %125 = OpIEqual %bool %124 %90
- OpSelectionMerge %126 None
- OpBranchConditional %125 %127 %126
- %127 = OpLabel
- %129 = OpLoad %v3float %color
- %130 = OpCompositeExtract %GammaTransferParams %params_0 3
- %128 = OpFunctionCall %v3float %gammaCorrection %129 %130
- OpStore %color %128
- %131 = OpCompositeExtract %mat3v3float %params_0 5
- %132 = OpLoad %v3float %color
- %133 = OpMatrixTimesVector %v3float %131 %132
- OpStore %color %133
- %135 = OpLoad %v3float %color
- %136 = OpCompositeExtract %GammaTransferParams %params_0 4
- %134 = OpFunctionCall %v3float %gammaCorrection %135 %136
- OpStore %color %134
- OpBranch %126
- %126 = OpLabel
- %137 = OpLoad %v3float %color
- %138 = OpCompositeExtract %float %137 0
- %139 = OpCompositeExtract %float %137 1
- %140 = OpCompositeExtract %float %137 2
- %141 = OpCompositeConstruct %v4float %138 %139 %140 %float_1
- OpReturnValue %141
+ %102 = OpCompositeExtract %uint %params_0 1
+ %104 = OpIEqual %bool %102 %103
+ OpSelectionMerge %105 None
+ OpBranchConditional %104 %106 %105
+ %106 = OpLabel
+ %108 = OpLoad %v3float %color
+ %109 = OpCompositeExtract %GammaTransferParams %params_0 3
+ %107 = OpFunctionCall %v3float %gammaCorrection %108 %109
+ OpStore %color %107
+ %110 = OpCompositeExtract %mat3v3float %params_0 5
+ %111 = OpLoad %v3float %color
+ %112 = OpMatrixTimesVector %v3float %110 %111
+ OpStore %color %112
+ %114 = OpLoad %v3float %color
+ %115 = OpCompositeExtract %GammaTransferParams %params_0 4
+ %113 = OpFunctionCall %v3float %gammaCorrection %114 %115
+ OpStore %color %113
+ OpBranch %105
+ %105 = OpLabel
+ %116 = OpLoad %v3float %color
+ %117 = OpCompositeExtract %float %116 0
+ %118 = OpCompositeExtract %float %116 1
+ %119 = OpCompositeExtract %float %116 2
+ %120 = OpCompositeConstruct %v4float %117 %118 %119 %float_1
+ OpReturnValue %120
OpFunctionEnd
-%conv_ExternalTextureParams = OpFunction %ExternalTextureParams None %142
+%conv_ExternalTextureParams = OpFunction %ExternalTextureParams None %121
%val = OpFunctionParameter %ExternalTextureParams_std140
- %145 = OpLabel
- %146 = OpCompositeExtract %uint %val 0
- %147 = OpCompositeExtract %uint %val 1
- %148 = OpCompositeExtract %mat3v4float %val 2
- %149 = OpCompositeExtract %GammaTransferParams %val 3
- %150 = OpCompositeExtract %GammaTransferParams %val 4
- %151 = OpCompositeExtract %mat3v3float %val 5
- %152 = OpCompositeExtract %v2float %val 6
- %153 = OpCompositeExtract %v2float %val 7
- %154 = OpCompositeExtract %v2float %val 8
- %155 = OpCompositeConstruct %mat3v2float %152 %153 %154
- %156 = OpCompositeConstruct %ExternalTextureParams %146 %147 %148 %149 %150 %151 %155
- OpReturnValue %156
+ %124 = OpLabel
+ %125 = OpCompositeExtract %uint %val 0
+ %126 = OpCompositeExtract %uint %val 1
+ %127 = OpCompositeExtract %mat3v4float %val 2
+ %128 = OpCompositeExtract %GammaTransferParams %val 3
+ %129 = OpCompositeExtract %GammaTransferParams %val 4
+ %130 = OpCompositeExtract %mat3v3float %val 5
+ %131 = OpCompositeExtract %v2float %val 6
+ %132 = OpCompositeExtract %v2float %val 7
+ %133 = OpCompositeExtract %v2float %val 8
+ %134 = OpCompositeConstruct %mat3v2float %131 %132 %133
+ %135 = OpCompositeConstruct %ExternalTextureParams %125 %126 %127 %128 %129 %130 %134
+ OpReturnValue %135
OpFunctionEnd
- %main = OpFunction %void None %157
- %160 = OpLabel
- %red = OpVariable %_ptr_Function_v4float Function %173
- %green = OpVariable %_ptr_Function_v4float Function %173
- %162 = OpLoad %3 %t
- %163 = OpLoad %3 %ext_tex_plane_1
- %169 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
- %170 = OpLoad %ExternalTextureParams_std140 %169
- %166 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %170
- %161 = OpFunctionCall %v4float %textureLoadExternal %162 %163 %165 %166
- OpStore %red %161
- %175 = OpLoad %19 %outImage
- %179 = OpLoad %19 %outImage
- %178 = OpImageQuerySize %v2uint %179
- %180 = OpISub %v2uint %178 %81
- %177 = OpBitcast %v2int %180
- %176 = OpFunctionCall %v2int %tint_clamp %95 %95 %177
- %181 = OpLoad %v4float %red
- OpImageWrite %175 %176 %181
- %183 = OpLoad %3 %t
- %184 = OpLoad %3 %ext_tex_plane_1
- %189 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
- %190 = OpLoad %ExternalTextureParams_std140 %189
- %188 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %190
- %182 = OpFunctionCall %v4float %textureLoadExternal %183 %184 %187 %188
- OpStore %green %182
- %193 = OpLoad %19 %outImage
- %200 = OpLoad %19 %outImage
- %199 = OpImageQuerySize %v2uint %200
- %201 = OpISub %v2uint %199 %81
- %198 = OpBitcast %v2int %201
- %194 = OpFunctionCall %v2int %tint_clamp %197 %95 %198
- %202 = OpLoad %v4float %green
- OpImageWrite %193 %194 %202
+ %main = OpFunction %void None %136
+ %139 = OpLabel
+ %red = OpVariable %_ptr_Function_v4float Function %159
+ %green = OpVariable %_ptr_Function_v4float Function %159
+ %141 = OpLoad %3 %t
+ %142 = OpLoad %3 %ext_tex_plane_1
+ %149 = OpLoad %3 %t
+ %148 = OpImageQuerySizeLod %v2uint %149 %int_0
+ %151 = OpISub %v2uint %148 %81
+ %147 = OpBitcast %v2int %151
+ %143 = OpFunctionCall %v2int %tint_clamp %145 %146 %147
+ %155 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
+ %156 = OpLoad %ExternalTextureParams_std140 %155
+ %152 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %156
+ %140 = OpFunctionCall %v4float %textureLoadExternal %141 %142 %143 %152
+ OpStore %red %140
+ %161 = OpLoad %19 %outImage
+ %165 = OpLoad %19 %outImage
+ %164 = OpImageQuerySize %v2uint %165
+ %166 = OpISub %v2uint %164 %81
+ %163 = OpBitcast %v2int %166
+ %162 = OpFunctionCall %v2int %tint_clamp %146 %146 %163
+ %167 = OpLoad %v4float %red
+ OpImageWrite %161 %162 %167
+ %169 = OpLoad %3 %t
+ %170 = OpLoad %3 %ext_tex_plane_1
+ %177 = OpLoad %3 %t
+ %176 = OpImageQuerySizeLod %v2uint %177 %int_0
+ %178 = OpISub %v2uint %176 %81
+ %175 = OpBitcast %v2int %178
+ %171 = OpFunctionCall %v2int %tint_clamp %174 %146 %175
+ %180 = OpAccessChain %_ptr_Uniform_ExternalTextureParams_std140 %ext_tex_params %uint_0
+ %181 = OpLoad %ExternalTextureParams_std140 %180
+ %179 = OpFunctionCall %ExternalTextureParams %conv_ExternalTextureParams %181
+ %168 = OpFunctionCall %v4float %textureLoadExternal %169 %170 %171 %179
+ OpStore %green %168
+ %184 = OpLoad %19 %outImage
+ %190 = OpLoad %19 %outImage
+ %189 = OpImageQuerySize %v2uint %190
+ %191 = OpISub %v2uint %189 %81
+ %188 = OpBitcast %v2int %191
+ %185 = OpFunctionCall %v2int %tint_clamp %187 %146 %188
+ %192 = OpLoad %v4float %green
+ OpImageWrite %184 %185 %192
OpReturn
OpFunctionEnd
diff --git a/test/tint/bug/tint/1739.wgsl.expected.wgsl b/test/tint/bug/tint/1739.wgsl.expected.wgsl
index 403d2d7..83cddcb 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.wgsl
@@ -1,60 +1,12 @@
-struct GammaTransferParams {
- G : f32,
- A : f32,
- B : f32,
- C : f32,
- D : f32,
- E : f32,
- F : f32,
- padding : u32,
-}
-
-struct ExternalTextureParams {
- numPlanes : u32,
- doYuvToRgbConversionOnly : u32,
- yuvToRgbConversionMatrix : mat3x4<f32>,
- gammaDecodeParams : GammaTransferParams,
- gammaEncodeParams : GammaTransferParams,
- gamutConversionMatrix : mat3x3<f32>,
- coordTransformationMatrix : mat3x2<f32>,
-}
-
-@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
-
-@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
-
-@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0) @binding(0) var t : texture_external;
@group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>;
-fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
- let cond = (abs(v) < vec3<f32>(params.D));
- let t = (sign(v) * ((params.C * abs(v)) + params.F));
- let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
- return select(f, t, cond);
-}
-
-fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
- let coord1 = (coord >> vec2<u32>(1));
- var color : vec3<f32>;
- if ((params.numPlanes == 1)) {
- color = textureLoad(plane0, coord, 0).rgb;
- } else {
- color = (vec4<f32>(textureLoad(plane0, coord, 0).r, textureLoad(plane1, coord1, 0).rg, 1) * params.yuvToRgbConversionMatrix);
- }
- if ((params.doYuvToRgbConversionOnly == 0)) {
- color = gammaCorrection(color, params.gammaDecodeParams);
- color = (params.gamutConversionMatrix * color);
- color = gammaCorrection(color, params.gammaEncodeParams);
- }
- return vec4<f32>(color, 1);
-}
-
@compute @workgroup_size(1)
fn main() {
- var red : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, vec2<i32>(10, 10), ext_tex_params);
+ var red : vec4<f32> = textureLoad(t, vec2<i32>(10, 10));
textureStore(outImage, vec2<i32>(0, 0), red);
- var green : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, vec2<i32>(70, 118), ext_tex_params);
+ var green : vec4<f32> = textureLoad(t, vec2<i32>(70, 118));
textureStore(outImage, vec2<i32>(1, 0), green);
return;
}
