diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 8d6a3b1dfa..c8d89e4420 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -906,8 +906,6 @@ libtint_source_set("libtint_writer_src") { "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 @@ if (tint_build_unittests) { "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 @@ if (tint_build_unittests) { ] } + # 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 @@ if (tint_build_unittests) { ":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 0f442ff720..0c2da7a7a2 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -550,8 +550,6 @@ list(APPEND TINT_LIB_SRCS 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 @@ if(TINT_BUILD_TESTS) 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 b4ff6d58ae..946bd36dc5 100644 --- a/src/tint/cmd/BUILD.gn +++ b/src/tint/cmd/BUILD.gn @@ -17,6 +17,8 @@ import("../../../tint_overrides_with_defaults.gni") 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 6da97019bc..142595b443 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 6e5daf0b9b..1b6ca74b6d 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 #include @@ -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 GenerateExternalTextureBindin } } - 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/writer/generate_external_texture_bindings.h b/src/tint/cmd/generate_external_texture_bindings.h similarity index 58% rename from src/tint/writer/generate_external_texture_bindings.h rename to src/tint/cmd/generate_external_texture_bindings.h index 8d0aad9e12..c6a842d4dc 100644 --- a/src/tint/writer/generate_external_texture_bindings.h +++ b/src/tint/cmd/generate_external_texture_bindings.h @@ -12,16 +12,15 @@ // 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_ +#ifndef SRC_TINT_CMD_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_ +#define SRC_TINT_CMD_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_ -#include "src/tint/transform/multiplanar_external_texture.h" +#include "tint/tint.h" -namespace tint::writer { +namespace tint::cmd { -transform::MultiplanarExternalTexture::BindingsMap GenerateExternalTextureBindings( - const Program* program); +writer::ExternalTextureOptions::BindingsMap GenerateExternalTextureBindings(const Program* program); -} // namespace tint::writer +} // namespace tint::cmd -#endif // SRC_TINT_WRITER_GENERATE_EXTERNAL_TEXTURE_BINDINGS_H_ +#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 b8625540e6..2c19675907 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 #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 @@ TEST_F(GenerateExternalTextureBindingsTest, Two_WithOtherBindingsInSameGroup) { } } // namespace -} // namespace tint::writer +} // namespace tint::cmd diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc index f4ad19fc8a..2daf75f22b 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 @@ bool GenerateSpirv(const tint::Program* program, const Options& options) { 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 @@ bool GenerateMsl(const tint::Program* program, const Options& options) { 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 @@ bool GenerateHlsl(const tint::Program* program, const Options& options) { 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 @@ 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.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 @@ int main(int argc, const char** argv) { m.Add(); 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 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(std::move(met_bindings)); - m.Add(); - return true; - }}, }; auto transform_names = [&] { tint::utils::StringStream names; diff --git a/src/tint/writer/glsl/generator.h b/src/tint/writer/glsl/generator.h index 5c2de66e68..ef7b65b450 100644 --- a/src/tint/writer/glsl/generator.h +++ b/src/tint/writer/glsl/generator.h @@ -74,9 +74,6 @@ 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 in the binding mappings for external textures ExternalTextureOptions external_texture_options = {}; @@ -87,7 +84,6 @@ struct Options { 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 29c28ff6f3..cdf4a15fea 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 @@ SanitizedResult Sanitize(const Program* in, manager.Add(); } - 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(new_bindings_map); + data.Add( + options.external_texture_options.bindings_map); manager.Add(); } diff --git a/src/tint/writer/hlsl/generator.h b/src/tint/writer/hlsl/generator.h index 1c072d8958..2cdfe6f799 100644 --- a/src/tint/writer/hlsl/generator.h +++ b/src/tint/writer/hlsl/generator.h @@ -59,9 +59,6 @@ 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 in the binding mappings for external textures ExternalTextureOptions external_texture_options = {}; @@ -80,7 +77,6 @@ struct Options { 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 84543ab65e..1cfbf64860 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 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { manager.Add(); } - 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(new_bindings_map); + data.Add( + options.external_texture_options.bindings_map); manager.Add(); } diff --git a/src/tint/writer/msl/generator.h b/src/tint/writer/msl/generator.h index e58ef9e49f..b4e11d9e23 100644 --- a/src/tint/writer/msl/generator.h +++ b/src/tint/writer/msl/generator.h @@ -63,9 +63,6 @@ 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 in the binding mappings for external textures ExternalTextureOptions external_texture_options = {}; @@ -79,7 +76,6 @@ struct Options { 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 397c96185f..268736f16f 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 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { manager.Add(); } - 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(new_bindings_map); + data.Add( + options.external_texture_options.bindings_map); manager.Add(); } diff --git a/src/tint/writer/spirv/generator.h b/src/tint/writer/spirv/generator.h index 11660fe6e5..5ad9057d79 100644 --- a/src/tint/writer/spirv/generator.h +++ b/src/tint/writer/spirv/generator.h @@ -46,9 +46,6 @@ 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 in the binding mappings for external textures ExternalTextureOptions external_texture_options = {}; @@ -60,7 +57,6 @@ struct Options { 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 3c74920a39..f3bbaf305d 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 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { manager.Add(); } - 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(new_bindings_map); + data.Add( + options.external_texture_options.bindings_map); manager.Add(); } diff --git a/test/tint/bug/tint/1739.wgsl b/test/tint/bug/tint/1739.wgsl index d513d64794..ef45dd5e1c 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; diff --git a/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl index e71fad6cc0..27ce96e902 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 @@ RWTexture2D outImage : register(u1, space0); 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 plane0, Texture2D 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 @@ ExternalTextureParams ext_tex_params_load(uint offset) { [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 e71fad6cc0..27ce96e902 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 @@ RWTexture2D outImage : register(u1, space0); 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 plane0, Texture2D 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 @@ ExternalTextureParams ext_tex_params_load(uint offset) { [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 3bbb84002c..c32e19b376 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(binding = 3, std140) uniform ext_tex_params_block_std140_ubo { 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 @@ vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ive 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 @@ void main() { 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 9bcbc79622..34134b318a 100644 --- a/test/tint/bug/tint/1739.wgsl.expected.msl +++ b/test/tint/bug/tint/1739.wgsl.expected.msl @@ -78,21 +78,18 @@ int2 tint_clamp(int2 e, int2 low, int2 high) { 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 plane0, texture2d 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 @@ float4 textureLoadExternal(texture2d plane0, texture2d tint_symbol_1 [[texture(0)]], texture2d tint_symbol_2 [[texture(1)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_3 [[buffer(0)]], texture2d 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 tint_symbol_1 [[texture(1)]], texture2d tint_symbol_2 [[texture(2)]], const constant ExternalTextureParams_tint_packed_vec3* tint_symbol_3 [[buffer(3)]], texture2d 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 b31df45344..01e7c869f0 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 403d2d7a74..83cddcbacb 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, - gammaDecodeParams : GammaTransferParams, - gammaEncodeParams : GammaTransferParams, - gamutConversionMatrix : mat3x3, - coordTransformationMatrix : mat3x2, -} - -@group(0) @binding(2) var ext_tex_plane_1 : texture_2d; - -@group(0) @binding(3) var ext_tex_params : ExternalTextureParams; - -@group(0) @binding(0) var t : texture_2d; +@group(0) @binding(0) var t : texture_external; @group(0) @binding(1) var outImage : texture_storage_2d; -fn gammaCorrection(v : vec3, params : GammaTransferParams) -> vec3 { - let cond = (abs(v) < vec3(params.D)); - let t = (sign(v) * ((params.C * abs(v)) + params.F)); - let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E)); - return select(f, t, cond); -} - -fn textureLoadExternal(plane0 : texture_2d, plane1 : texture_2d, coord : vec2, params : ExternalTextureParams) -> vec4 { - let coord1 = (coord >> vec2(1)); - var color : vec3; - if ((params.numPlanes == 1)) { - color = textureLoad(plane0, coord, 0).rgb; - } else { - color = (vec4(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(color, 1); -} - @compute @workgroup_size(1) fn main() { - var red : vec4 = textureLoadExternal(t, ext_tex_plane_1, vec2(10, 10), ext_tex_params); + var red : vec4 = textureLoad(t, vec2(10, 10)); textureStore(outImage, vec2(0, 0), red); - var green : vec4 = textureLoadExternal(t, ext_tex_plane_1, vec2(70, 118), ext_tex_params); + var green : vec4 = textureLoad(t, vec2(70, 118)); textureStore(outImage, vec2(1, 0), green); return; }