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>
This commit is contained in:
dan sinclair
Dawn LUCI CQ
parent
a56069b03d
commit
6521a357f5
|
|
@ -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",
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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",
|
||||
]
|
||||
|
|
|
|||
|
|
@ -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"
|
||||
|
|
|
|||
|
|
@ -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 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
|
||||
|
|
@ -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_
|
||||
|
|
@ -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 @@ TEST_F(GenerateExternalTextureBindingsTest, Two_WithOtherBindingsInSameGroup) {
|
|||
}
|
||||
|
||||
} // namespace
|
||||
} // namespace tint::writer
|
||||
} // namespace tint::cmd
|
||||
|
|
@ -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<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;
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
};
|
||||
|
|
|
|||
|
|
@ -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<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>();
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
};
|
||||
|
|
|
|||
|
|
@ -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<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>();
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
};
|
||||
|
|
|
|||
|
|
@ -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<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>();
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
};
|
||||
|
|
|
|||
|
|
@ -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<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>();
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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>;
|
||||
|
|
|
|||
|
|
@ -31,33 +31,18 @@ RWTexture2D<float4> 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<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 @@ 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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -31,33 +31,18 @@ RWTexture2D<float4> 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<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 @@ 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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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.
|
||||
|
||||
|
||||
|
||||
|
|
|
|||
|
|
@ -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<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 @@ float4 textureLoadExternal(texture2d<float, access::sample> plane0, texture2d<fl
|
|||
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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue