Add TextureBuiltinsFromUniform transform
TextureNumLevels/TextureNumSamples counterpart builtin functions
are not available in GLSL for OpenGLES.
This transform replace those builtin calls by creating an internal
uniform buffer to store these builtin values. A bindpoint to
byte offset and data type is included in the result of the transform
so that caller will know how to upload the data.
But: dawn:1299, tint:2006
Change-Id: Idc1b6e699443b8b656208d49e73ec6001f1e77ea
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/145621
Commit-Queue: Shrek Shao <shrekshao@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/include/tint/tint.h b/include/tint/tint.h
index c8f6581..899de9a 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -25,6 +25,7 @@
#include "src/tint/api/options/array_length_from_uniform.h"
#include "src/tint/api/options/binding_remapper.h"
#include "src/tint/api/options/external_texture.h"
+#include "src/tint/api/options/texture_builtins_from_uniform.h"
#include "src/tint/api/tint.h"
#include "src/tint/lang/core/type/manager.h"
#include "src/tint/lang/wgsl/ast/transform/first_index_offset.h"
diff --git a/src/tint/api/options/BUILD.cmake b/src/tint/api/options/BUILD.cmake
index 2edad38..8117549 100644
--- a/src/tint/api/options/BUILD.cmake
+++ b/src/tint/api/options/BUILD.cmake
@@ -29,6 +29,7 @@
api/options/binding_remapper.h
api/options/external_texture.h
api/options/options.cc
+ api/options/texture_builtins_from_uniform.h
)
tint_target_add_dependencies(tint_api_options lib
diff --git a/src/tint/api/options/BUILD.gn b/src/tint/api/options/BUILD.gn
index e24f6f1..9608ec4 100644
--- a/src/tint/api/options/BUILD.gn
+++ b/src/tint/api/options/BUILD.gn
@@ -30,6 +30,7 @@
"binding_remapper.h",
"external_texture.h",
"options.cc",
+ "texture_builtins_from_uniform.h",
]
deps = [
"${tint_src_dir}/api/common",
diff --git a/src/tint/api/options/texture_builtins_from_uniform.h b/src/tint/api/options/texture_builtins_from_uniform.h
new file mode 100644
index 0000000..4dc2eff
--- /dev/null
+++ b/src/tint/api/options/texture_builtins_from_uniform.h
@@ -0,0 +1,52 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_API_OPTIONS_TEXTURE_BUILTINS_FROM_UNIFORM_H_
+#define SRC_TINT_API_OPTIONS_TEXTURE_BUILTINS_FROM_UNIFORM_H_
+
+#include <unordered_map>
+#include <utility>
+
+#include "src/tint/api/common/binding_point.h"
+#include "src/tint/utils/reflection/reflection.h"
+
+namespace tint {
+
+/// Options used to specify a mapping of binding points to indices into a UBO
+/// from which to load buffer sizes.
+struct TextureBuiltinsFromUniformOptions {
+ /// Indicate the type of field for each entry to push.
+ enum class Field {
+ /// The number of mip levels of the bonnd texture view.
+ TextureNumLevels,
+ /// The number of samples per texel of the bound multipsampled texture.
+ TextureNumSamples,
+ };
+
+ /// Records the field and the byte offset of the data to push in the internal uniform buffer.
+ using FieldAndOffset = std::pair<Field, uint32_t>;
+ /// Maps from binding point to data entry with the information to populate the data.
+ using BindingPointToFieldAndOffset = std::unordered_map<BindingPoint, FieldAndOffset>;
+
+ /// The binding point to use to generate a uniform buffer from which to read
+ /// buffer sizes.
+ BindingPoint ubo_binding = {};
+
+ /// Reflect the fields of this class so that it can be used by tint::ForeachField()
+ TINT_REFLECT(ubo_binding);
+};
+
+} // namespace tint
+
+#endif // SRC_TINT_API_OPTIONS_TEXTURE_BUILTINS_FROM_UNIFORM_H_
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 127a2da..895fd56 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -62,6 +62,7 @@
#include "src/tint/lang/wgsl/ast/transform/single_entry_point.h"
#include "src/tint/lang/wgsl/ast/transform/std140.h"
#include "src/tint/lang/wgsl/ast/transform/texture_1d_to_2d.h"
+#include "src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h"
#include "src/tint/lang/wgsl/ast/transform/unshadow.h"
#include "src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.h"
#include "src/tint/lang/wgsl/ast/variable_decl_statement.h"
@@ -229,6 +230,15 @@
manager.Add<ast::transform::RemovePhonies>();
+ // TextureBuiltinsFromUniform must come before CombineSamplers to preserve texture binding point
+ // info, instead of combined sampler binding point. As a result, TextureBuiltinsFromUniform also
+ // comes before BindingRemapper so the binding point info it reflects is before remapping.
+ if (options.texture_builtins_from_uniform) {
+ manager.Add<ast::transform::TextureBuiltinsFromUniform>();
+ data.Add<ast::transform::TextureBuiltinsFromUniform::Config>(
+ options.texture_builtins_from_uniform->ubo_binding);
+ }
+
data.Add<ast::transform::CombineSamplers::BindingInfo>(options.binding_map,
options.placeholder_binding_point);
manager.Add<ast::transform::CombineSamplers>();
@@ -254,6 +264,10 @@
SanitizedResult result;
ast::transform::DataMap outputs;
result.program = manager.Run(in, data, outputs);
+ if (auto* res = outputs.Get<ast::transform::TextureBuiltinsFromUniform::Result>()) {
+ result.needs_internal_uniform_buffer = true;
+ result.bindpoint_to_data = std::move(res->bindpoint_to_data);
+ }
return result;
}
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
index 6c8f5fb..ad44120 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
@@ -21,6 +21,7 @@
#include <unordered_set>
#include <utility>
+#include "src/tint/api/options/texture_builtins_from_uniform.h"
#include "src/tint/lang/core/builtin_value.h"
#include "src/tint/lang/glsl/writer/common/version.h"
#include "src/tint/lang/wgsl/program/program_builder.h"
@@ -53,6 +54,13 @@
/// The sanitized program.
Program program;
+
+ /// True if the shader needs a UBO.
+ bool needs_internal_uniform_buffer = false;
+
+ /// Store a map of global texture variable binding point to the byte offset and data type to
+ /// push into the internal uniform buffer.
+ TextureBuiltinsFromUniformOptions::BindingPointToFieldAndOffset bindpoint_to_data;
};
/// Sanitize a program in preparation for generating GLSL.
diff --git a/src/tint/lang/glsl/writer/common/options.h b/src/tint/lang/glsl/writer/common/options.h
index e2aadd7..2f55c4b 100644
--- a/src/tint/lang/glsl/writer/common/options.h
+++ b/src/tint/lang/glsl/writer/common/options.h
@@ -15,10 +15,12 @@
#ifndef SRC_TINT_LANG_GLSL_WRITER_COMMON_OPTIONS_H_
#define SRC_TINT_LANG_GLSL_WRITER_COMMON_OPTIONS_H_
+#include <optional>
#include <string>
#include <unordered_map>
#include "src/tint/api/options/external_texture.h"
+#include "src/tint/api/options/texture_builtins_from_uniform.h"
#include "src/tint/lang/core/access.h"
#include "src/tint/lang/glsl/writer/common/version.h"
#include "src/tint/lang/wgsl/sem/sampler_texture_pair.h"
@@ -66,6 +68,11 @@
/// Options used in the binding mappings for external textures
ExternalTextureOptions external_texture_options = {};
+ /// Options used to map WGSL textureNumLevels/textureNumSamples builtins to internal uniform
+ /// buffer values. If not specified, emits corresponding GLSL builtins
+ /// textureQueryLevels/textureSamples directly.
+ std::optional<TextureBuiltinsFromUniformOptions> texture_builtins_from_uniform = std::nullopt;
+
/// The GLSL version to emit
Version version;
@@ -74,6 +81,7 @@
allow_collisions,
disable_workgroup_init,
external_texture_options,
+ texture_builtins_from_uniform,
version);
};
diff --git a/src/tint/lang/glsl/writer/output.h b/src/tint/lang/glsl/writer/output.h
index 084b0ef..286af3f 100644
--- a/src/tint/lang/glsl/writer/output.h
+++ b/src/tint/lang/glsl/writer/output.h
@@ -19,6 +19,7 @@
#include <utility>
#include <vector>
+#include "src/tint/api/options/texture_builtins_from_uniform.h"
#include "src/tint/lang/wgsl/ast/pipeline_stage.h"
namespace tint::glsl::writer {
@@ -39,6 +40,13 @@
/// The list of entry points in the generated GLSL.
std::vector<std::pair<std::string, ast::PipelineStage>> entry_points;
+
+ /// True if the shader needs a UBO.
+ bool needs_internal_uniform_buffer = false;
+
+ /// Store a map of global texture variable binding points to the byte offset and data type to
+ /// push into the internal uniform buffer.
+ TextureBuiltinsFromUniformOptions::BindingPointToFieldAndOffset bindpoint_to_data;
};
} // namespace tint::glsl::writer
diff --git a/src/tint/lang/glsl/writer/writer.cc b/src/tint/lang/glsl/writer/writer.cc
index 4124346..e02c6bb 100644
--- a/src/tint/lang/glsl/writer/writer.cc
+++ b/src/tint/lang/glsl/writer/writer.cc
@@ -15,6 +15,7 @@
#include "src/tint/lang/glsl/writer/writer.h"
#include <memory>
+#include <utility>
#include "src/tint/lang/glsl/writer/ast_printer/ast_printer.h"
#include "src/tint/lang/wgsl/ast/transform/binding_remapper.h"
@@ -43,6 +44,8 @@
Output output;
output.glsl = impl->Result();
+ output.needs_internal_uniform_buffer = sanitized_result.needs_internal_uniform_buffer;
+ output.bindpoint_to_data = std::move(sanitized_result.bindpoint_to_data);
// Collect the list of entry points in the sanitized program.
for (auto* func : sanitized_result.program.AST().Functions()) {
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.cmake b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
index e56ea18..d91332e 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.cmake
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
@@ -115,6 +115,8 @@
lang/wgsl/ast/transform/substitute_override.h
lang/wgsl/ast/transform/texture_1d_to_2d.cc
lang/wgsl/ast/transform/texture_1d_to_2d.h
+ lang/wgsl/ast/transform/texture_builtins_from_uniform.cc
+ lang/wgsl/ast/transform/texture_builtins_from_uniform.h
lang/wgsl/ast/transform/transform.cc
lang/wgsl/ast/transform/transform.h
lang/wgsl/ast/transform/truncate_interstage_variables.cc
@@ -213,6 +215,7 @@
lang/wgsl/ast/transform/std140_test.cc
lang/wgsl/ast/transform/substitute_override_test.cc
lang/wgsl/ast/transform/texture_1d_to_2d_test.cc
+ lang/wgsl/ast/transform/texture_builtins_from_uniform_test.cc
lang/wgsl/ast/transform/transform_test.cc
lang/wgsl/ast/transform/truncate_interstage_variables_test.cc
lang/wgsl/ast/transform/unshadow_test.cc
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.gn b/src/tint/lang/wgsl/ast/transform/BUILD.gn
index 8df2f94..c55205d 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.gn
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.gn
@@ -120,6 +120,8 @@
"substitute_override.h",
"texture_1d_to_2d.cc",
"texture_1d_to_2d.h",
+ "texture_builtins_from_uniform.cc",
+ "texture_builtins_from_uniform.h",
"transform.cc",
"transform.h",
"truncate_interstage_variables.cc",
@@ -216,6 +218,7 @@
"std140_test.cc",
"substitute_override_test.cc",
"texture_1d_to_2d_test.cc",
+ "texture_builtins_from_uniform_test.cc",
"transform_test.cc",
"truncate_interstage_variables_test.cc",
"unshadow_test.cc",
diff --git a/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.cc b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.cc
new file mode 100644
index 0000000..4c4baee
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.cc
@@ -0,0 +1,491 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h"
+
+#include <memory>
+#include <queue>
+#include <string>
+#include <utility>
+#include <variant>
+#include <vector>
+
+#include "src/tint/lang/wgsl/program/clone_context.h"
+#include "src/tint/lang/wgsl/program/program_builder.h"
+#include "src/tint/lang/wgsl/resolver/resolve.h"
+#include "src/tint/lang/wgsl/sem/call.h"
+#include "src/tint/lang/wgsl/sem/function.h"
+#include "src/tint/lang/wgsl/sem/module.h"
+#include "src/tint/lang/wgsl/sem/statement.h"
+#include "src/tint/lang/wgsl/sem/variable.h"
+
+#include "src/tint/utils/containers/hashmap.h"
+#include "src/tint/utils/containers/vector.h"
+#include "src/tint/utils/rtti/switch.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::TextureBuiltinsFromUniform);
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::TextureBuiltinsFromUniform::Config);
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::TextureBuiltinsFromUniform::Result);
+
+namespace tint::ast::transform {
+
+namespace {
+
+/// The member name of the texture builtin values.
+constexpr std::string_view kTextureBuiltinValuesMemberNamePrefix = "texture_builtin_value_";
+
+bool ShouldRun(const Program* program) {
+ for (auto* fn : program->AST().Functions()) {
+ if (auto* sem_fn = program->Sem().Get(fn)) {
+ for (auto* builtin : sem_fn->DirectlyCalledBuiltins()) {
+ // GLSL ES has no native support for the counterpart of
+ // textureNumLevels (textureQueryLevels) and textureNumSamples (textureSamples)
+ if (builtin->Type() == core::Function::kTextureNumLevels) {
+ return true;
+ }
+ if (builtin->Type() == core::Function::kTextureNumSamples) {
+ return true;
+ }
+ }
+ }
+ }
+ return false;
+}
+
+} // namespace
+
+TextureBuiltinsFromUniform::TextureBuiltinsFromUniform() = default;
+TextureBuiltinsFromUniform::~TextureBuiltinsFromUniform() = default;
+
+/// PIMPL state for the transform
+struct TextureBuiltinsFromUniform::State {
+ /// Constructor
+ /// @param program the source program
+ /// @param in the input transform data
+ /// @param out the output transform data
+ explicit State(const Program* program, const DataMap& in, DataMap& out)
+ : src(program), inputs(in), outputs(out) {}
+
+ /// Runs the transform
+ /// @returns the new program or SkipTransform if the transform is not required
+ ApplyResult Run() {
+ auto* cfg = inputs.Get<Config>();
+ if (cfg == nullptr) {
+ b.Diagnostics().add_error(
+ diag::System::Transform,
+ "missing transform data for " +
+ std::string(tint::TypeInfo::Of<TextureBuiltinsFromUniform>().name));
+ return resolver::Resolve(b);
+ }
+
+ if (!ShouldRun(ctx.src)) {
+ return SkipTransform;
+ }
+
+ // The dependency order declartions guaranteed that we traverse interested functions in the
+ // following order:
+ // 1. texture builtins
+ // 2. user function directly calls texture builtins
+ // 3. user function calls 2.
+ // 4. user function calls 3.
+ // ...
+ // n. entry point function.
+ for (auto* fn_decl : sem.Module()->DependencyOrderedDeclarations()) {
+ if (auto* fn = sem.Get<sem::Function>(fn_decl)) {
+ for (auto* call : fn->DirectCalls()) {
+ auto* call_expr = call->Declaration();
+
+ tint::Switch(
+ call->Target(),
+ [&](const sem::Builtin* builtin) {
+ if (builtin->Type() != core::Function::kTextureNumLevels &&
+ builtin->Type() != core::Function::kTextureNumSamples) {
+ return;
+ }
+ if (auto* call_stmt =
+ call->Stmt()->Declaration()->As<CallStatement>()) {
+ if (call_stmt->expr == call->Declaration()) {
+ // textureNumLevels() / textureNumSamples() is used as a
+ // statement. The argument expression must be side-effect free,
+ // so just drop the statement.
+ RemoveStatement(ctx, call_stmt);
+ return;
+ }
+ }
+
+ auto* texture_expr = call->Declaration()->args[0];
+ auto* texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
+ TINT_ASSERT(texture_sem);
+
+ TextureBuiltinsFromUniformOptions::Field dataType =
+ GetFieldFromBuiltinFunctionType(builtin->Type());
+
+ tint::Switch(
+ texture_sem,
+ [&](const sem::GlobalVariable* global) {
+ // This texture variable is a global variable.
+ auto binding = GetAndRecordGlobalBinding(global, dataType);
+ // Record the call and binding to be replaced later.
+ builtin_to_replace.Add(call_expr, binding);
+ },
+ [&](const sem::Variable* variable) {
+ // This texture variable is a user function parameter.
+ auto new_param =
+ GetAndRecordFunctionParameter(fn, variable, dataType);
+ // Record the call and new_param to be replaced later.
+ builtin_to_replace.Add(call_expr, new_param);
+ },
+ [&](Default) {
+ TINT_ICE() << "unexpected texture root identifier";
+ });
+ },
+ [&](const sem::Function* user_fn) {
+ auto user_param_to_info = fn_to_data.Find(user_fn);
+ if (!user_param_to_info) {
+ // Uninterested function not calling texture builtins with function
+ // texture param.
+ return;
+ }
+ TINT_ASSERT(call->Arguments().Length() ==
+ user_fn->Declaration()->params.Length());
+ for (size_t i = 0; i < call->Arguments().Length(); i++) {
+ auto param = user_fn->Declaration()->params[i];
+ auto info = user_param_to_info->Get(param);
+ if (info.has_value()) {
+ auto* arg = call->Arguments()[i];
+ auto* texture_sem = arg->RootIdentifier();
+ auto& args = call_to_data.GetOrCreate(call_expr, [&] {
+ return Vector<
+ std::variant<BindingPoint, const ast::Parameter*>, 4>();
+ });
+
+ tint::Switch(
+ texture_sem,
+ [&](const sem::GlobalVariable* global) {
+ // This texture variable is a global variable.
+ auto binding =
+ GetAndRecordGlobalBinding(global, info->field);
+ // Record the binding to add to args.
+ args.Push(binding);
+ },
+ [&](const sem::Variable* variable) {
+ // This texture variable is a user function parameter.
+ auto new_param = GetAndRecordFunctionParameter(
+ fn, variable, info->field);
+ // Record adding extra function parameter
+ args.Push(new_param);
+ },
+ [&](Default) {
+ TINT_ICE() << "unexpected texture root identifier";
+ });
+ }
+ }
+ });
+ }
+ }
+ }
+
+ // If there's no interested texture builtin at all, skip the transform.
+ if (bindpoint_to_data.empty()) {
+ return SkipTransform;
+ }
+
+ // If any functions need extra params, add them now.
+ if (!fn_to_data.IsEmpty()) {
+ for (auto pair : fn_to_data) {
+ auto* fn = pair.key;
+
+ // Reorder the param to a vector to make sure params are in the correct order.
+ Vector<const ast::Parameter*, 4> extra_params_in_order;
+ extra_params_in_order.Resize(pair.value.Count());
+ for (auto t_p : pair.value) {
+ TINT_ASSERT(t_p.value.extra_idx < extra_params_in_order.Length());
+ extra_params_in_order[t_p.value.extra_idx] = t_p.value.param;
+ }
+
+ for (auto p : extra_params_in_order) {
+ ctx.InsertBack(fn->Declaration()->params, p);
+ }
+ }
+ }
+
+ // Replace all interested texture builtin calls.
+ for (auto pair : builtin_to_replace) {
+ auto call = pair.key;
+ if (std::holds_alternative<BindingPoint>(pair.value)) {
+ // This texture is a global variable with binding point.
+ // Read builtin value from uniform buffer.
+ auto* builtin_value = GetUniformValue(std::get<BindingPoint>(pair.value));
+ ctx.Replace(call, builtin_value);
+ } else {
+ // Otherwise this value comes from a function param
+ auto* param = std::get<const ast::Parameter*>(pair.value);
+ ctx.Replace(call, b.Expr(param));
+ }
+ }
+
+ // Insert all extra args to interested function calls.
+ for (auto pair : call_to_data) {
+ auto call = pair.key;
+ for (auto new_arg_info : pair.value) {
+ if (std::holds_alternative<BindingPoint>(new_arg_info)) {
+ // This texture is a global variable with binding point.
+ // Read builtin value from uniform buffer.
+ auto* builtin_value = GetUniformValue(std::get<BindingPoint>(new_arg_info));
+ ctx.InsertBack(call->args, builtin_value);
+ } else {
+ // Otherwise this value comes from a function param
+ auto* param = std::get<const ast::Parameter*>(new_arg_info);
+ ctx.InsertBack(call->args, b.Expr(param));
+ }
+ }
+ }
+
+ outputs.Add<Result>(bindpoint_to_data);
+
+ ctx.Clone();
+ return resolver::Resolve(b);
+ }
+
+ private:
+ /// The source program
+ const Program* const src;
+ /// The transform inputs
+ const DataMap& inputs;
+ /// The transform outputs
+ DataMap& outputs;
+ /// The target program builder
+ ProgramBuilder b;
+ /// The clone context
+ program::CloneContext ctx = {&b, src, /* auto_clone_symbols */ true};
+ /// Alias to the semantic info in ctx.src
+ const sem::Info& sem = ctx.src->Sem();
+
+ /// The bindpoint to byte offset and field to pass out in transform result.
+ /// For one texture type, it could only be passed into one of the
+ /// textureNumLevels or textureNumSamples because their accepting param texture
+ /// type is different. There cannot be a binding entry with both field type.
+ /// Note: because this transform must be run before CombineSampler and BindingRemapper,
+ /// the binding number here is before remapped.
+ Result::BindingPointToFieldAndOffset bindpoint_to_data;
+
+ struct FunctionExtraParamInfo {
+ using Field = TextureBuiltinsFromUniformOptions::Field;
+ // The kind of texture information this parameter holds.
+ Field field = Field::TextureNumLevels;
+
+ // The extra passed in param that corresponds to the texture param.
+ const ast::Parameter* param = nullptr;
+
+ // id of this extra param e.g. f(t0, foo, t1, e0, e1) e0 and e1 are extra params, their
+ // extra_idx are 0 and 1. This is to help sort extra ids in the correct order.
+ size_t extra_idx = 0;
+ };
+
+ /// Store a map from function to a collection of extra params that need adding.
+ /// The value of the map is made a map instead of a vector to make it easier to find the param.
+ /// for call sites. e.g. fn f(t: texture_2d<f32>) -> u32 {
+ /// return textureNumLevels(t);
+ /// }
+ /// ->
+ /// fn f(t : texture_2d<f32>, tint_symbol : u32) -> u32 {
+ /// return tint_symbol;
+ /// }
+ Hashmap<const sem::Function*, Hashmap<const ast::Parameter*, FunctionExtraParamInfo, 4>, 8>
+ fn_to_data;
+
+ /// For each callsite of the above functions, record a vector of extra call args that need
+ /// inserting. e.g. f(tex)
+ /// ->
+ /// f(tex, internal_uniform.texture_builtin_value), if tex is from a global
+ /// variable, store the BindingPoint. or f(tex, extra_param_tex), if tex is from a function
+ /// param, store the texture function parameter pointer.
+ Hashmap<const CallExpression*, Vector<std::variant<BindingPoint, const ast::Parameter*>, 4>, 8>
+ call_to_data;
+
+ /// Texture builtin calls to be replaced by either uniform values or function parameters.
+ Hashmap<const CallExpression*, std::variant<BindingPoint, const ast::Parameter*>, 8>
+ builtin_to_replace;
+
+ /// A map from global texture bindpoint to the symbol storing its builtin value in the uniform
+ /// buffer struct.
+ Hashmap<BindingPoint, Symbol, 16> bindpoint_to_syms;
+
+ /// The internal uniform buffer
+ const Variable* ubo = nullptr;
+ /// Get or create a UBO including u32 scalars for texture builtin values.
+ /// @returns the symbol of the uniform buffer variable.
+ Symbol GetUboSym() {
+ if (ubo) {
+ // Already created
+ return ubo->name->symbol;
+ }
+
+ auto* cfg = inputs.Get<Config>();
+
+ Vector<const ast::StructMember*, 16> new_members;
+ new_members.Resize(bindpoint_to_data.size());
+ for (auto it : bindpoint_to_data) {
+ // Emit a u32 scalar for each binding that needs builtin value passed in.
+ size_t i = it.second.second / sizeof(uint32_t);
+ TINT_ASSERT(i < new_members.Length());
+ // Append the vector index with the variable name to avoid unstable naming issue.
+ auto sym = b.Symbols().New(std::string(kTextureBuiltinValuesMemberNamePrefix) +
+ std::to_string(i));
+ bindpoint_to_syms.Add(it.first, sym);
+ new_members[i] = b.Member(sym, b.ty.u32());
+ }
+
+ // Find if there's any existing global variable using the same cfg->ubo_binding
+ for (auto* var : src->AST().Globals<Var>()) {
+ if (var->HasBindingPoint()) {
+ auto* global_sem = sem.Get<sem::GlobalVariable>(var);
+
+ // The original binding point
+ BindingPoint binding_point = *global_sem->BindingPoint();
+
+ if (binding_point == cfg->ubo_binding) {
+ // This ubo_binding struct already exists.
+ // which should only be added by other *FromUniform transforms.
+ // Replace it with a new struct including the new_member.
+ // Then remove the old structure global declaration.
+
+ ubo = var->As<Variable>();
+
+ auto* ty = global_sem->Type()->UnwrapRef();
+ auto* str = ty->As<sem::Struct>();
+ if (TINT_UNLIKELY(!str)) {
+ TINT_ICE()
+ << "existing ubo binding " << cfg->ubo_binding << " is not a struct.";
+ return ctx.Clone(ubo->name->symbol);
+ }
+
+ for (auto new_member : new_members) {
+ ctx.InsertBack(str->Declaration()->members, new_member);
+ }
+ return ctx.Clone(ubo->name->symbol);
+ }
+ }
+ }
+
+ auto* buffer_struct = b.Structure(b.Sym(), std::move(new_members));
+ ubo = b.GlobalVar(b.Sym(), b.ty.Of(buffer_struct), core::AddressSpace::kUniform,
+ b.Group(core::AInt(cfg->ubo_binding.group)),
+ b.Binding(core::AInt(cfg->ubo_binding.binding)));
+ return ubo->name->symbol;
+ }
+
+ /// Get the expression of retrieving the builtin value from the uniform buffer.
+ /// @param binding of the global variable.
+ /// @returns an expression of the builtin value.
+ const ast::Expression* GetUniformValue(const BindingPoint& binding) {
+ auto iter = bindpoint_to_data.find(binding);
+ TINT_ASSERT(iter != bindpoint_to_data.end());
+
+ // Make sure GetUboSym() is called first to initialize the uniform buffer struct.
+ auto ubo_sym = GetUboSym();
+ // Load the builtin value from the UBO.
+ auto member_sym = bindpoint_to_syms.Get(binding);
+ TINT_ASSERT(member_sym.has_value());
+ auto* builtin_value = b.MemberAccessor(ubo_sym, *member_sym);
+
+ return builtin_value;
+ }
+
+ /// Get and return the binding of the global texture variable. Record in bindpoint_to_data if
+ /// first visited.
+ /// @param global global variable of the texture variable.
+ /// @param field type of the interested builtin function data related to this texture.
+ /// @returns binding of the global variable.
+ BindingPoint GetAndRecordGlobalBinding(const sem::GlobalVariable* global,
+ TextureBuiltinsFromUniformOptions::Field field) {
+ auto binding = global->BindingPoint().value();
+ auto iter = bindpoint_to_data.find(binding);
+ if (iter == bindpoint_to_data.end()) {
+ // First visit, recording the binding.
+ uint32_t index = static_cast<uint32_t>(bindpoint_to_data.size());
+ bindpoint_to_data.emplace(
+ binding,
+ Result::FieldAndOffset{field, index * static_cast<uint32_t>(sizeof(uint32_t))});
+ }
+ return binding;
+ }
+
+ /// Find which function param is the given texture variable.
+ /// Add a new u32 param relates to this texture param. Record in fn_to_data if first visited.
+ /// @param fn the current function scope.
+ /// @param var the texture variable.
+ /// @param field type of the interested builtin function data related to this texture.
+ /// @returns the new u32 function parameter.
+ const ast::Parameter* GetAndRecordFunctionParameter(
+ const sem::Function* fn,
+ const sem::Variable* var,
+ TextureBuiltinsFromUniformOptions::Field field) {
+ auto& param_to_info = fn_to_data.GetOrCreate(
+ fn, [&] { return Hashmap<const ast::Parameter*, FunctionExtraParamInfo, 4>(); });
+
+ const ast::Parameter* param = nullptr;
+ for (auto p : fn->Declaration()->params) {
+ if (p->As<Variable>() == var->Declaration()) {
+ param = p;
+ break;
+ }
+ }
+ TINT_ASSERT(param);
+ // Get or record a new u32 param to this function if first visited.
+ auto entry = param_to_info.Get(param);
+ if (entry.has_value()) {
+ return entry->param;
+ }
+ const ast::Parameter* new_param = b.Param(b.Sym(), b.ty.u32());
+ size_t idx = param_to_info.Count();
+ param_to_info.Add(param, FunctionExtraParamInfo{field, new_param, idx});
+ return new_param;
+ }
+
+ /// Get the uniform options field for the builtin function.
+ /// @param type of the builtin function
+ /// @returns corresponding TextureBuiltinsFromUniformOptions::Field for the builtin
+ static TextureBuiltinsFromUniformOptions::Field GetFieldFromBuiltinFunctionType(
+ core::Function type) {
+ switch (type) {
+ case core::Function::kTextureNumLevels:
+ return TextureBuiltinsFromUniformOptions::Field::TextureNumLevels;
+ case core::Function::kTextureNumSamples:
+ return TextureBuiltinsFromUniformOptions::Field::TextureNumSamples;
+ default:
+ TINT_UNREACHABLE() << "unsupported builtin function type " << type;
+ }
+ return TextureBuiltinsFromUniformOptions::Field::TextureNumLevels;
+ }
+};
+
+Transform::ApplyResult TextureBuiltinsFromUniform::Apply(const Program* src,
+ const DataMap& inputs,
+ DataMap& outputs) const {
+ return State{src, inputs, outputs}.Run();
+}
+
+TextureBuiltinsFromUniform::Config::Config(BindingPoint ubo_bp) : ubo_binding(ubo_bp) {}
+TextureBuiltinsFromUniform::Config::Config(const Config&) = default;
+TextureBuiltinsFromUniform::Config& TextureBuiltinsFromUniform::Config::operator=(const Config&) =
+ default;
+TextureBuiltinsFromUniform::Config::~Config() = default;
+
+TextureBuiltinsFromUniform::Result::Result(BindingPointToFieldAndOffset bindpoint_to_data_in)
+ : bindpoint_to_data(std::move(bindpoint_to_data_in)) {}
+TextureBuiltinsFromUniform::Result::Result(const Result&) = default;
+TextureBuiltinsFromUniform::Result::~Result() = default;
+
+} // namespace tint::ast::transform
diff --git a/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h
new file mode 100644
index 0000000..0db043e
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h
@@ -0,0 +1,119 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_LANG_WGSL_AST_TRANSFORM_TEXTURE_BUILTINS_FROM_UNIFORM_H_
+#define SRC_TINT_LANG_WGSL_AST_TRANSFORM_TEXTURE_BUILTINS_FROM_UNIFORM_H_
+
+#include <unordered_map>
+#include <unordered_set>
+
+#include "src/tint/api/common/binding_point.h"
+#include "src/tint/api/options/texture_builtins_from_uniform.h"
+#include "src/tint/lang/wgsl/ast/transform/transform.h"
+
+// Forward declarations
+namespace tint {
+class CloneContext;
+} // namespace tint
+
+namespace tint::ast::transform {
+
+/// TextureBuiltinsFromUniform is a transform that implements calls to textureNumLevels() and
+/// textureNumSamples() by retrieving the texture information from a uniform buffer, as those
+/// builtin functions are not available in some version of GLSL.
+///
+/// The generated uniform buffer will have the form:
+/// ```
+/// struct internal_uniform {
+/// texture_builtin_value_0 : u32,
+/// };
+///
+/// @group(0) @binding(0) var tex : texture_2d<f32>;
+/// ```
+/// The binding group and number used for this uniform buffer are provided via
+/// the `Config` transform input.
+///
+/// The transform coverts the texture builtins calls into values lookup from the internal
+/// buffer. If the texture is a function parameter instead of a global variable, this transform
+/// also takes care of adding extra paramters and arguments to these functions and their callsites.
+///
+/// This transform must run before `CombineSamplers` transform so that the binding point of the
+/// original texture object can be preserved.
+class TextureBuiltinsFromUniform final : public Castable<TextureBuiltinsFromUniform, Transform> {
+ public:
+ /// Constructor
+ TextureBuiltinsFromUniform();
+ /// Destructor
+ ~TextureBuiltinsFromUniform() override;
+
+ /// Configuration options for the TextureBuiltinsFromUniform transform.
+ struct Config final : public Castable<Config, Data> {
+ /// Constructor
+ /// @param ubo_bp the binding point to use for the generated uniform buffer.
+ explicit Config(BindingPoint ubo_bp);
+
+ /// Copy constructor
+ Config(const Config&);
+
+ /// Copy assignment
+ /// @return this Config
+ Config& operator=(const Config&);
+
+ /// Destructor
+ ~Config() override;
+
+ /// The binding point to use for the generated uniform buffer.
+ BindingPoint ubo_binding;
+ };
+
+ /// Information produced about what the transform did.
+ /// If there were no calls to the textureNumLevels() or textureNumSamples() builtin, then no
+ /// Result will be emitted.
+ struct Result final : public Castable<Result, Data> {
+ /// Using for shorter names
+ /// Records the field and the byte offset of the data to push in the internal uniform
+ /// buffer.
+ using FieldAndOffset = TextureBuiltinsFromUniformOptions::FieldAndOffset;
+ /// Maps from binding point to data entry with the information to populate the data.
+ using BindingPointToFieldAndOffset =
+ TextureBuiltinsFromUniformOptions::BindingPointToFieldAndOffset;
+
+ /// Constructor
+ /// @param bindpoint_to_data_in mapping from binding points of global texture variables to
+ /// the byte offsets and data types needed to be pushed into the internal uniform buffer.
+ explicit Result(BindingPointToFieldAndOffset bindpoint_to_data_in);
+
+ /// Copy constructor
+ Result(const Result&);
+
+ /// Destructor
+ ~Result() override;
+
+ /// A map of global texture variable binding point to the byte offset and data type to push
+ /// into the internal uniform buffer.
+ BindingPointToFieldAndOffset bindpoint_to_data;
+ };
+
+ /// @copydoc Transform::Apply
+ ApplyResult Apply(const Program* program,
+ const DataMap& inputs,
+ DataMap& outputs) const override;
+
+ private:
+ struct State;
+};
+
+} // namespace tint::ast::transform
+
+#endif // SRC_TINT_LANG_WGSL_AST_TRANSFORM_TEXTURE_BUILTINS_FROM_UNIFORM_H_
diff --git a/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform_test.cc b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform_test.cc
new file mode 100644
index 0000000..21185e8
--- /dev/null
+++ b/src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform_test.cc
@@ -0,0 +1,687 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/lang/wgsl/ast/transform/texture_builtins_from_uniform.h"
+
+#include <utility>
+
+#include "src/tint/lang/wgsl/ast/transform/helper_test.h"
+
+namespace tint::ast::transform {
+namespace {
+
+using TextureBuiltinsFromUniformTest = TransformTest;
+
+TEST_F(TextureBuiltinsFromUniformTest, ShouldRunEmptyModule) {
+ auto* src = R"()";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ EXPECT_FALSE(ShouldRun<TextureBuiltinsFromUniform>(src, data));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, ShouldRunNoTextureNumLevels) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ _ = textureDimensions(t);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ EXPECT_FALSE(ShouldRun<TextureBuiltinsFromUniform>(src, data));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, ShouldRunWithTextureNumLevels) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumLevels(t);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ EXPECT_TRUE(ShouldRun<TextureBuiltinsFromUniform>(src, data));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, Error_MissingTransformData) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumLevels(t);
+}
+)";
+
+ auto* expect =
+ "error: missing transform data for tint::ast::transform::TextureBuiltinsFromUniform";
+
+ auto got = Run<TextureBuiltinsFromUniform>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, BasicTextureNumLevels) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumLevels(t);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = tint_symbol_1.texture_builtin_value_0;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ // Note: Using the following EXPECT_EQ directly on BindingPointToFieldAndOffset seems to cause
+ // compiler to hang. EXPECT_EQ(
+ // TextureBuiltinsFromUniformOptions::BindingPointToFieldAndOffset{
+ // {BindgPoint{0u, 0u},
+ // std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u)}},
+ // val->bindpoint_to_data);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, BasicTextureNumSamples) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var samples : u32 = textureNumSamples(t);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var t : texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var samples : u32 = tint_symbol_1.texture_builtin_value_0;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumSamples, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, SameBuiltinCalledMultipleTimes) {
+ auto* src = R"(
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumLevels(tex);
+ len = textureNumLevels(tex);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = tint_symbol_1.texture_builtin_value_0;
+ len = tint_symbol_1.texture_builtin_value_0;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, SameBuiltinCalledMultipleTimesTextureNumSamples) {
+ auto* src = R"(
+@group(0) @binding(0) var tex : texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumSamples(tex);
+ len = textureNumSamples(tex);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var tex : texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = tint_symbol_1.texture_builtin_value_0;
+ len = tint_symbol_1.texture_builtin_value_0;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumSamples, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, TextureAsFunctionParameterBasic) {
+ auto* src = R"(
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f(t: texture_2d<f32>) -> u32 {
+ return textureNumLevels(t);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_1 {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_2 : tint_symbol_1;
+
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f(t : texture_2d<f32>, tint_symbol : u32) -> u32 {
+ return tint_symbol;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex, tint_symbol_2.texture_builtin_value_0);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, TextureAsFunctionParameterUsedTwice) {
+ auto* src = R"(
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f(t: texture_2d<f32>) -> u32 {
+ var len = textureNumLevels(t);
+ len += textureNumLevels(t);
+ return len;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_1 {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_2 : tint_symbol_1;
+
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f(t : texture_2d<f32>, tint_symbol : u32) -> u32 {
+ var len = tint_symbol;
+ len += tint_symbol;
+ return len;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex, tint_symbol_2.texture_builtin_value_0);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, TextureAsFunctionParameterMultipleParameters) {
+ auto* src = R"(
+@group(0) @binding(0) var tex1 : texture_2d<f32>;
+@group(0) @binding(1) var tex2 : texture_2d<f32>;
+@group(0) @binding(2) var tex3 : texture_2d<f32>;
+
+fn f(t1: texture_2d<f32>, t2: texture_2d<f32>, t3: texture_2d<f32>) -> u32 {
+ return textureNumLevels(t1) + textureNumLevels(t2) + textureNumLevels(t3);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex1, tex2, tex3);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_3 {
+ texture_builtin_value_0 : u32,
+ texture_builtin_value_1 : u32,
+ texture_builtin_value_2 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_4 : tint_symbol_3;
+
+@group(0) @binding(0) var tex1 : texture_2d<f32>;
+
+@group(0) @binding(1) var tex2 : texture_2d<f32>;
+
+@group(0) @binding(2) var tex3 : texture_2d<f32>;
+
+fn f(t1 : texture_2d<f32>, t2 : texture_2d<f32>, t3 : texture_2d<f32>, tint_symbol : u32, tint_symbol_1 : u32, tint_symbol_2 : u32) -> u32 {
+ return ((tint_symbol + tint_symbol_1) + tint_symbol_2);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f(tex1, tex2, tex3, tint_symbol_4.texture_builtin_value_0, tint_symbol_4.texture_builtin_value_1, tint_symbol_4.texture_builtin_value_2);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(3u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 4u),
+ val->bindpoint_to_data.at(BindingPoint{0, 1}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 8u),
+ val->bindpoint_to_data.at(BindingPoint{0, 2}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, TextureAsFunctionParameterNested) {
+ auto* src = R"(
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f2(tt: texture_2d<f32>) -> u32 {
+ return textureNumLevels(tt);
+}
+
+fn f1(t: texture_2d<f32>) -> u32 {
+ return f2(t);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f1(tex);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_2 {
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+@group(0) @binding(0) var tex : texture_2d<f32>;
+
+fn f2(tt : texture_2d<f32>, tint_symbol : u32) -> u32 {
+ return tint_symbol;
+}
+
+fn f1(t : texture_2d<f32>, tint_symbol_1 : u32) -> u32 {
+ return f2(t, tint_symbol_1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = f1(tex, tint_symbol_3.texture_builtin_value_0);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(1u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, TextureAsFunctionParameterMixed) {
+ auto* src = R"(
+@group(0) @binding(0) var tex0 : texture_2d<f32>;
+@group(0) @binding(1) var tex1 : texture_2d<f32>;
+@group(0) @binding(2) var tex2 : texture_2d<f32>;
+@group(0) @binding(3) var tex3 : texture_2d<f32>;
+@group(0) @binding(4) var tex4 : texture_2d_array<f32>; // unused for textureNumLevels
+
+fn f_nested(t1: texture_2d<f32>, t2: texture_2d<f32>) -> u32 {
+ return textureNumLevels(t1) + textureNumLevels(t2);
+}
+
+fn f1(a: u32, t: texture_2d<f32>) -> u32 {
+ return a + f_nested(t, tex1) + textureNumLevels(tex3);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ _ = textureNumLayers(tex4);
+ _ = f1(9u, tex0);
+ _ = f_nested(tex2, tex2);
+ _ = f_nested(tex1, tex0);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_3 {
+ texture_builtin_value_0 : u32,
+ texture_builtin_value_1 : u32,
+ texture_builtin_value_2 : u32,
+ texture_builtin_value_3 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_4 : tint_symbol_3;
+
+@group(0) @binding(0) var tex0 : texture_2d<f32>;
+
+@group(0) @binding(1) var tex1 : texture_2d<f32>;
+
+@group(0) @binding(2) var tex2 : texture_2d<f32>;
+
+@group(0) @binding(3) var tex3 : texture_2d<f32>;
+
+@group(0) @binding(4) var tex4 : texture_2d_array<f32>;
+
+fn f_nested(t1 : texture_2d<f32>, t2 : texture_2d<f32>, tint_symbol : u32, tint_symbol_1 : u32) -> u32 {
+ return (tint_symbol + tint_symbol_1);
+}
+
+fn f1(a : u32, t : texture_2d<f32>, tint_symbol_2 : u32) -> u32 {
+ return ((a + f_nested(t, tex1, tint_symbol_2, tint_symbol_4.texture_builtin_value_0)) + tint_symbol_4.texture_builtin_value_1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ _ = textureNumLayers(tex4);
+ _ = f1(9u, tex0, tint_symbol_4.texture_builtin_value_2);
+ _ = f_nested(tex2, tex2, tint_symbol_4.texture_builtin_value_3, tint_symbol_4.texture_builtin_value_3);
+ _ = f_nested(tex1, tex0, tint_symbol_4.texture_builtin_value_0, tint_symbol_4.texture_builtin_value_2);
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(4u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 1}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 4u),
+ val->bindpoint_to_data.at(BindingPoint{0, 3}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 8u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 12u),
+ val->bindpoint_to_data.at(BindingPoint{0, 2}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, MultipleTextures) {
+ auto* src = R"(
+@group(0) @binding(0) var t0 : texture_2d<f32>;
+@group(0) @binding(1) var t1 : texture_multisampled_2d<f32>;
+@group(0) @binding(2) var t2 : texture_2d_array<f32>;
+@group(0) @binding(3) var t3 : texture_cube<f32>;
+@group(0) @binding(4) var t4 : texture_depth_2d;
+@group(1) @binding(0) var t5 : texture_depth_multisampled_2d;
+
+@compute @workgroup_size(1)
+fn main() {
+ _ = textureNumLevels(t0);
+ _ = textureNumSamples(t1);
+ _ = textureNumLevels(t2);
+ _ = textureNumLevels(t3);
+ _ = textureNumLevels(t4);
+ _ = textureNumSamples(t5);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ texture_builtin_value_0 : u32,
+ texture_builtin_value_1 : u32,
+ texture_builtin_value_2 : u32,
+ texture_builtin_value_3 : u32,
+ texture_builtin_value_4 : u32,
+ texture_builtin_value_5 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var t0 : texture_2d<f32>;
+
+@group(0) @binding(1) var t1 : texture_multisampled_2d<f32>;
+
+@group(0) @binding(2) var t2 : texture_2d_array<f32>;
+
+@group(0) @binding(3) var t3 : texture_cube<f32>;
+
+@group(0) @binding(4) var t4 : texture_depth_2d;
+
+@group(1) @binding(0) var t5 : texture_depth_multisampled_2d;
+
+@compute @workgroup_size(1)
+fn main() {
+ _ = tint_symbol_1.texture_builtin_value_0;
+ _ = tint_symbol_1.texture_builtin_value_1;
+ _ = tint_symbol_1.texture_builtin_value_2;
+ _ = tint_symbol_1.texture_builtin_value_3;
+ _ = tint_symbol_1.texture_builtin_value_4;
+ _ = tint_symbol_1.texture_builtin_value_5;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(6u, val->bindpoint_to_data.size());
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumSamples, 4u),
+ val->bindpoint_to_data.at(BindingPoint{0, 1}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 8u),
+ val->bindpoint_to_data.at(BindingPoint{0, 2}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 12u),
+ val->bindpoint_to_data.at(BindingPoint{0, 3}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 16u),
+ val->bindpoint_to_data.at(BindingPoint{0, 4}));
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumSamples, 20u),
+ val->bindpoint_to_data.at(BindingPoint{1, 0}));
+}
+
+TEST_F(TextureBuiltinsFromUniformTest, BindingPointExist) {
+ auto* src = R"(
+struct tint_symbol {
+ foo : array<vec4<u32>, 1u>,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = textureNumLevels(t);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ foo : array<vec4<u32>, 1u>,
+ texture_builtin_value_0 : u32,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = tint_symbol_1.texture_builtin_value_0;
+}
+)";
+
+ TextureBuiltinsFromUniform::Config cfg({0, 30u});
+
+ DataMap data;
+ data.Add<TextureBuiltinsFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<TextureBuiltinsFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ auto* val = got.data.Get<TextureBuiltinsFromUniform::Result>();
+ ASSERT_NE(val, nullptr);
+ EXPECT_EQ(std::make_pair(TextureBuiltinsFromUniformOptions::Field::TextureNumLevels, 0u),
+ val->bindpoint_to_data.at(BindingPoint{0, 0}));
+}
+
+} // namespace
+} // namespace tint::ast::transform
