[tint][ast] Reimplement ArrayLengthFromUniform transform to support ptr params
Bug: tint:2053
Change-Id: I33e43aa2a9b4bb6f040166e0a3bcedf5b84f17be
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/176462
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/hlsl/writer/ast_printer/sanitizer_test.cc b/src/tint/lang/hlsl/writer/ast_printer/sanitizer_test.cc
index 5823771..72bfc3b 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/sanitizer_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/sanitizer_test.cc
@@ -166,17 +166,17 @@
ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
auto got = gen.Result();
- auto* expect = R"(cbuffer cbuffer_tint_symbol_1 : register(b4, space3) {
- uint4 tint_symbol_1[2];
+ auto* expect = R"(cbuffer cbuffer_tint_array_lengths : register(b4, space3) {
+ uint4 tint_array_lengths[2];
};
ByteAddressBuffer b : register(t1, space2);
ByteAddressBuffer c : register(t2, space2);
void a_func() {
- uint tint_symbol_3 = 0u;
- b.GetDimensions(tint_symbol_3);
- uint tint_symbol_4 = ((tint_symbol_3 - 0u) / 4u);
- uint len = (tint_symbol_4 + ((tint_symbol_1[1].w - 0u) / 4u));
+ uint tint_symbol_1 = 0u;
+ b.GetDimensions(tint_symbol_1);
+ uint tint_symbol_2 = ((tint_symbol_1 - 0u) / 4u);
+ uint len = (tint_symbol_2 + ((tint_array_lengths[1].w - 0u) / 4u));
return;
}
)";
diff --git a/src/tint/lang/msl/writer/ast_printer/sanitizer_test.cc b/src/tint/lang/msl/writer/ast_printer/sanitizer_test.cc
index 5bebc55..901ddd2 100644
--- a/src/tint/lang/msl/writer/ast_printer/sanitizer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/sanitizer_test.cc
@@ -78,16 +78,16 @@
T elements[N];
};
-struct tint_symbol {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct my_struct {
tint_array<float, 1> a;
};
-fragment void a_func(const constant tint_symbol* tint_symbol_2 [[buffer(30)]]) {
- uint len = (((*(tint_symbol_2)).buffer_size[0u][1u] - 0u) / 4u);
+fragment void a_func(const constant TintArrayLengths* tint_symbol [[buffer(30)]]) {
+ uint len = (((*(tint_symbol)).array_lengths[0u][1u] - 0u) / 4u);
return;
}
@@ -135,8 +135,8 @@
T elements[N];
};
-struct tint_symbol {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct my_struct {
@@ -144,8 +144,8 @@
tint_array<float, 1> a;
};
-fragment void a_func(const constant tint_symbol* tint_symbol_2 [[buffer(30)]]) {
- uint len = (((*(tint_symbol_2)).buffer_size[0u][1u] - 4u) / 4u);
+fragment void a_func(const constant TintArrayLengths* tint_symbol [[buffer(30)]]) {
+ uint len = (((*(tint_symbol)).array_lengths[0u][1u] - 4u) / 4u);
return;
}
@@ -196,16 +196,16 @@
T elements[N];
};
-struct tint_symbol {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct my_struct {
tint_array<float, 1> a;
};
-fragment void a_func(const constant tint_symbol* tint_symbol_2 [[buffer(30)]]) {
- uint len = (((*(tint_symbol_2)).buffer_size[0u][1u] - 0u) / 4u);
+fragment void a_func(const constant TintArrayLengths* tint_symbol [[buffer(30)]]) {
+ uint len = (((*(tint_symbol)).array_lengths[0u][1u] - 0u) / 4u);
return;
}
@@ -256,16 +256,16 @@
T elements[N];
};
-struct tint_symbol {
- /* 0x0000 */ tint_array<uint4, 2> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 2> array_lengths;
};
struct my_struct {
tint_array<float, 1> a;
};
-fragment void a_func(const constant tint_symbol* tint_symbol_2 [[buffer(29)]]) {
- uint len = ((((*(tint_symbol_2)).buffer_size[1u][3u] - 0u) / 4u) + (((*(tint_symbol_2)).buffer_size[0u][2u] - 0u) / 4u));
+fragment void a_func(const constant TintArrayLengths* tint_symbol [[buffer(29)]]) {
+ uint len = ((((*(tint_symbol)).array_lengths[1u][3u] - 0u) / 4u) + (((*(tint_symbol)).array_lengths[0u][2u] - 0u) / 4u));
return;
}
diff --git a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
index bdb43cf..e307e25 100644
--- a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
+++ b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
@@ -29,24 +29,39 @@
#include <memory>
#include <string>
+#include <string_view>
#include <utility>
#include "src/tint/lang/core/fluent_types.h"
+#include "src/tint/lang/core/unary_op.h"
+#include "src/tint/lang/wgsl/ast/expression.h"
#include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
+#include "src/tint/lang/wgsl/ast/unary_op_expression.h"
+#include "src/tint/lang/wgsl/ast/variable.h"
+#include "src/tint/lang/wgsl/builtin_fn.h"
#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/array.h"
+#include "src/tint/lang/wgsl/sem/builtin_fn.h"
#include "src/tint/lang/wgsl/sem/call.h"
+#include "src/tint/lang/wgsl/sem/expression.h"
#include "src/tint/lang/wgsl/sem/function.h"
+#include "src/tint/lang/wgsl/sem/member_accessor_expression.h"
#include "src/tint/lang/wgsl/sem/statement.h"
#include "src/tint/lang/wgsl/sem/variable.h"
+#include "src/tint/utils/containers/unique_vector.h"
+#include "src/tint/utils/diagnostic/diagnostic.h"
+#include "src/tint/utils/ice/ice.h"
+#include "src/tint/utils/rtti/switch.h"
+#include "src/tint/utils/text/text_style.h"
TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::ArrayLengthFromUniform);
TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::ArrayLengthFromUniform::Config);
TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::ArrayLengthFromUniform::Result);
using namespace tint::core::fluent_types; // NOLINT
- //
+
namespace tint::ast::transform {
namespace {
@@ -75,12 +90,11 @@
/// @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) {}
+ : src(program), outputs(out), cfg(in.Get<Config>()) {}
/// 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().AddError(diag::System::Transform, Source{})
<< "missing transform data for "
@@ -88,103 +102,34 @@
return resolver::Resolve(b);
}
- if (!ShouldRun(src)) {
+ if (cfg->bindpoint_to_size_index.empty() || !ShouldRun(src)) {
return SkipTransform;
}
- const char* kBufferSizeMemberName = "buffer_size";
+ // Create the name of the array lengths uniform variable.
+ array_lengths_var = b.Symbols().New("tint_array_lengths");
- // Determine the size of the buffer size array.
- uint32_t max_buffer_size_index = 0;
-
- IterateArrayLengthOnStorageVar(
- [&](const CallExpression*, const sem::VariableUser*, const sem::GlobalVariable* var) {
- if (auto binding = var->Attributes().binding_point) {
- auto idx_itr = cfg->bindpoint_to_size_index.find(*binding);
- if (idx_itr == cfg->bindpoint_to_size_index.end()) {
- return;
- }
- if (idx_itr->second > max_buffer_size_index) {
- max_buffer_size_index = idx_itr->second;
+ // Replace all the arrayLength() calls.
+ for (auto* fn : src.AST().Functions()) {
+ if (auto* sem_fn = sem.Get(fn)) {
+ for (auto* call : sem_fn->DirectCalls()) {
+ if (auto* target = call->Target()->As<sem::BuiltinFn>()) {
+ if (target->Fn() == wgsl::BuiltinFn::kArrayLength) {
+ ReplaceArrayLengthCall(call);
+ }
}
}
- });
-
- // Get (or create, on first call) the uniform buffer that will receive the
- // size of each storage buffer in the module.
- const Variable* buffer_size_ubo = nullptr;
- auto get_ubo = [&] {
- if (!buffer_size_ubo) {
- // Emit an array<vec4<u32>, N>, where N is 1/4 number of elements.
- // We do this because UBOs require an element stride that is 16-byte
- // aligned.
- auto* buffer_size_struct = b.Structure(
- b.Sym(), tint::Vector{
- b.Member(kBufferSizeMemberName,
- b.ty.array(b.ty.vec4(b.ty.u32()),
- u32((max_buffer_size_index / 4) + 1))),
- });
- buffer_size_ubo =
- b.GlobalVar(b.Sym(), b.ty.Of(buffer_size_struct), core::AddressSpace::kUniform,
- b.Group(AInt(cfg->ubo_binding.group)),
- b.Binding(AInt(cfg->ubo_binding.binding)));
}
- return buffer_size_ubo;
- };
+ }
- std::unordered_set<uint32_t> used_size_indices;
+ // Add the necessary array-length arguments to all the newly created array-length
+ // parameters.
+ while (!len_params_needing_args.IsEmpty()) {
+ AddArrayLengthArguments(len_params_needing_args.Pop());
+ }
- IterateArrayLengthOnStorageVar([&](const CallExpression* call_expr,
- const sem::VariableUser* storage_buffer_sem,
- const sem::GlobalVariable* var) {
- auto binding = var->Attributes().binding_point;
- if (!binding) {
- return;
- }
- auto idx_itr = cfg->bindpoint_to_size_index.find(*binding);
- if (idx_itr == cfg->bindpoint_to_size_index.end()) {
- return;
- }
-
- uint32_t size_index = idx_itr->second;
- used_size_indices.insert(size_index);
-
- // Load the total storage buffer size from the UBO.
- uint32_t array_index = size_index / 4;
- auto* vec_expr = b.IndexAccessor(
- b.MemberAccessor(get_ubo()->name->symbol, kBufferSizeMemberName), u32(array_index));
- uint32_t vec_index = size_index % 4;
- auto* total_storage_buffer_size = b.IndexAccessor(vec_expr, u32(vec_index));
-
- // Calculate actual array length
- // total_storage_buffer_size - array_offset
- // array_length = ----------------------------------------
- // array_stride
- const Expression* total_size = total_storage_buffer_size;
- if (TINT_UNLIKELY(storage_buffer_sem->Type()->Is<core::type::Pointer>())) {
- TINT_ICE() << "storage buffer variable should not be a pointer. These should have "
- "been removed by the SimplifyPointers transform";
- return;
- }
- auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
- const core::type::Array* array_type = nullptr;
- if (auto* str = storage_buffer_type->As<core::type::Struct>()) {
- // The variable is a struct, so subtract the byte offset of the array
- // member.
- auto* array_member_sem = str->Members().Back();
- array_type = array_member_sem->Type()->As<core::type::Array>();
- total_size = b.Sub(total_storage_buffer_size, u32(array_member_sem->Offset()));
- } else if (auto* arr = storage_buffer_type->As<core::type::Array>()) {
- array_type = arr;
- } else {
- TINT_ICE() << "expected form of arrayLength argument to be &array_var or "
- "&struct_var.array_member";
- return;
- }
- auto* array_length = b.Div(total_size, u32(array_type->Stride()));
-
- ctx.Replace(call_expr, array_length);
- });
+ // Add the tint_array_lengths module-scope uniform variable.
+ AddArrayLengthsUniformVar();
outputs.Add<Result>(used_size_indices);
@@ -193,81 +138,192 @@
}
private:
+ // Replaces the arrayLength() builtin call with an array-length expression passed via a uniform
+ // buffer.
+ void ReplaceArrayLengthCall(const sem::Call* call) {
+ if (auto* replacement = ArrayLengthOf(call->Arguments()[0])) {
+ ctx.Replace(call->Declaration(), replacement);
+ }
+ }
+
+ /// @returns an AST expression that is equal to the arrayLength() of the runtime-sized array
+ /// accessed by the pointer expression @p expr, or nullptr on error or if the array is not in
+ /// the Config::bindpoint_to_size_index map.
+ const ast::Expression* ArrayLengthOf(const sem::Expression* expr) {
+ const ast::Expression* len = nullptr;
+ while (expr) {
+ expr = Switch(
+ expr, //
+ [&](const sem::VariableUser* user) {
+ len = ArrayLengthOf(user->Variable());
+ return nullptr;
+ },
+ [&](const sem::MemberAccessorExpression* access) {
+ return access->Object(); // Follow the object
+ },
+ [&](const sem::Expression* e) {
+ return Switch(
+ e->Declaration(), //
+ [&](const ast::UnaryOpExpression* unary) -> const sem::Expression* {
+ switch (unary->op) {
+ case core::UnaryOp::kAddressOf:
+ case core::UnaryOp::kIndirection:
+ return sem.Get(unary->expr); // Follow the object
+ default:
+ TINT_ICE() << "unexpected unary op: " << unary->op;
+ return nullptr;
+ }
+ },
+ TINT_ICE_ON_NO_MATCH);
+ },
+ TINT_ICE_ON_NO_MATCH);
+ }
+ return len;
+ }
+
+ /// @returns an AST expression that is equal to the arrayLength() of the runtime-sized array
+ /// held by the module-scope variable or parameter @p var, or nullptr on error or if the array
+ /// is not in the Config::bindpoint_to_size_index map.
+ const ast::Expression* ArrayLengthOf(const sem::Variable* var) {
+ return Switch(
+ var, //
+ [&](const sem::GlobalVariable* global) { return ArrayLengthOf(global); },
+ [&](const sem::Parameter* param) { return ArrayLengthOf(param); },
+ TINT_ICE_ON_NO_MATCH);
+ }
+
+ /// @returns an AST expression that is equal to the arrayLength() of the runtime-sized array
+ /// held by the module scope variable @p global, or nullptr on error or if the array is not in
+ /// the Config::bindpoint_to_size_index map.
+ const ast::Expression* ArrayLengthOf(const sem::GlobalVariable* global) {
+ auto binding = global->Attributes().binding_point;
+ TINT_ASSERT_OR_RETURN_VALUE(binding, nullptr);
+
+ auto idx_it = cfg->bindpoint_to_size_index.find(*binding);
+ if (idx_it == cfg->bindpoint_to_size_index.end()) {
+ // If the bindpoint_to_size_index map does not contain an entry for the storage buffer,
+ // then we preserve the arrayLength() call.
+ return nullptr;
+ }
+
+ uint32_t size_index = idx_it->second;
+ used_size_indices.insert(size_index);
+
+ // Load the total storage buffer size from the UBO.
+ uint32_t array_index = size_index / 4;
+ auto* vec_expr = b.IndexAccessor(
+ b.MemberAccessor(array_lengths_var, kArrayLengthsMemberName), u32(array_index));
+ uint32_t vec_index = size_index % 4;
+ auto* total_storage_buffer_size = b.IndexAccessor(vec_expr, u32(vec_index));
+
+ // Calculate actual array length
+ // total_storage_buffer_size - array_offset
+ // array_length = ----------------------------------------
+ // array_stride
+ const Expression* total_size = total_storage_buffer_size;
+ if (TINT_UNLIKELY(global->Type()->Is<core::type::Pointer>())) {
+ TINT_ICE() << "storage buffer variable should not be a pointer. "
+ "These should have been removed by the SimplifyPointers transform";
+ return nullptr;
+ }
+ auto* storage_buffer_type = global->Type()->UnwrapRef();
+ const core::type::Array* array_type = nullptr;
+ if (auto* str = storage_buffer_type->As<core::type::Struct>()) {
+ // The variable is a struct, so subtract the byte offset of the
+ // array member.
+ auto* array_member_sem = str->Members().Back();
+ array_type = array_member_sem->Type()->As<core::type::Array>();
+ total_size = b.Sub(total_storage_buffer_size, u32(array_member_sem->Offset()));
+ } else if (auto* arr = storage_buffer_type->As<core::type::Array>()) {
+ array_type = arr;
+ } else {
+ TINT_ICE() << "expected form of arrayLength argument to be &array_var or "
+ "&struct_var.array_member";
+ return nullptr;
+ }
+ return b.Div(total_size, u32(array_type->Stride()));
+ }
+
+ /// @returns an AST expression that is equal to the arrayLength() of the runtime-sized array
+ /// held by the object pointed to by the pointer parameter @p param.
+ const ast::Expression* ArrayLengthOf(const sem::Parameter* param) {
+ // Pointer originates from a parameter.
+ // Add a new array length parameter to the function, and use that.
+ auto len_name = param_lengths.GetOrAdd(param, [&] {
+ auto* fn = param->Owner()->As<sem::Function>();
+ auto name = b.Symbols().New(param->Declaration()->name->symbol.Name() + "_length");
+ auto* len_param = b.Param(name, b.ty.u32());
+ ctx.InsertAfter(fn->Declaration()->params, param->Declaration(), len_param);
+ len_params_needing_args.Add(param);
+ return name;
+ });
+ return b.Expr(len_name);
+ }
+
+ /// Constructs the uniform buffer variable that will hold the array lengths.
+ void AddArrayLengthsUniformVar() {
+ // Calculate the highest index in the array lengths array
+ uint32_t highest_index = 0;
+ for (auto idx : used_size_indices) {
+ if (idx > highest_index) {
+ highest_index = idx;
+ }
+ }
+
+ // Emit an array<vec4<u32>, N>, where N is 1/4 number of elements.
+ // We do this because UBOs require an element stride that is 16-byte aligned.
+ auto* buffer_size_struct =
+ b.Structure(b.Symbols().New("TintArrayLengths"),
+ tint::Vector{
+ b.Member(kArrayLengthsMemberName,
+ b.ty.array(b.ty.vec4<u32>(), u32((highest_index / 4) + 1))),
+ });
+ b.GlobalVar(array_lengths_var, b.ty.Of(buffer_size_struct), core::AddressSpace::kUniform,
+ b.Group(AInt(cfg->ubo_binding.group)),
+ b.Binding(AInt(cfg->ubo_binding.binding)));
+ }
+
+ /// Adds an additional array-length argument to all the calls to the function that owns the
+ /// pointer parameter @p param. This may add new entries to #len_params_needing_args.
+ void AddArrayLengthArguments(const sem::Parameter* param) {
+ auto* fn = param->Owner()->As<sem::Function>();
+ for (auto* call : fn->CallSites()) {
+ auto* arg = call->Arguments()[param->Index()];
+ if (auto* len = ArrayLengthOf(arg); len) {
+ ctx.InsertAfter(call->Declaration()->args, arg->Declaration(), len);
+ } else {
+ // Callee expects an array length, but there's no binding for it.
+ // Call arrayLength() at the call-site.
+ len = b.Call(wgsl::BuiltinFn::kArrayLength, ctx.Clone(arg->Declaration()));
+ ctx.InsertAfter(call->Declaration()->args, arg->Declaration(), len);
+ }
+ }
+ }
+
+ /// Name of the array-lengths struct member that holds all the array lengths.
+ static constexpr std::string_view kArrayLengthsMemberName = "array_lengths";
+
/// The source program
const Program& src;
- /// The transform inputs
- const DataMap& inputs;
/// The transform outputs
DataMap& outputs;
+ /// The transform config
+ const Config* const cfg;
/// The target program builder
ProgramBuilder b;
/// The clone context
program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true};
-
- /// Iterate over all arrayLength() builtins that operate on
- /// storage buffer variables.
- /// @param functor of type void(const CallExpression*, const
- /// sem::VariableUser, const sem::GlobalVariable*). It takes in an
- /// CallExpression of the arrayLength call expression node, a
- /// sem::VariableUser of the used storage buffer variable, and the
- /// sem::GlobalVariable for the storage buffer.
- template <typename F>
- void IterateArrayLengthOnStorageVar(F&& functor) {
- auto& sem = src.Sem();
-
- // Find all calls to the arrayLength() builtin.
- for (auto* node : src.ASTNodes().Objects()) {
- auto* call_expr = node->As<CallExpression>();
- if (!call_expr) {
- continue;
- }
-
- auto* call = sem.Get(call_expr)->UnwrapMaterialize()->As<sem::Call>();
- auto* builtin = call->Target()->As<sem::BuiltinFn>();
- if (!builtin || builtin->Fn() != wgsl::BuiltinFn::kArrayLength) {
- continue;
- }
-
- if (auto* call_stmt = call->Stmt()->Declaration()->As<CallStatement>()) {
- if (call_stmt->expr == call_expr) {
- // arrayLength() is used as a statement.
- // The argument expression must be side-effect free, so just drop the statement.
- RemoveStatement(ctx, call_stmt);
- continue;
- }
- }
-
- // Get the storage buffer that contains the runtime array.
- // Since we require SimplifyPointers, we can assume that the arrayLength()
- // call has one of two forms:
- // arrayLength(&struct_var.array_member)
- // arrayLength(&array_var)
- auto* param = call_expr->args[0]->As<UnaryOpExpression>();
- if (TINT_UNLIKELY(!param || param->op != core::UnaryOp::kAddressOf)) {
- TINT_ICE() << "expected form of arrayLength argument to be &array_var or "
- "&struct_var.array_member";
- break;
- }
- auto* storage_buffer_expr = param->expr;
- if (auto* accessor = param->expr->As<MemberAccessorExpression>()) {
- storage_buffer_expr = accessor->object;
- }
- auto* storage_buffer_sem = sem.Get<sem::VariableUser>(storage_buffer_expr);
- if (TINT_UNLIKELY(!storage_buffer_sem)) {
- TINT_ICE() << "expected form of arrayLength argument to be &array_var or "
- "&struct_var.array_member";
- break;
- }
-
- // Get the index to use for the buffer size array.
- auto* var = tint::As<sem::GlobalVariable>(storage_buffer_sem->Variable());
- if (TINT_UNLIKELY(!var)) {
- TINT_ICE() << "storage buffer is not a global variable";
- break;
- }
- functor(call_expr, storage_buffer_sem, var);
- }
- }
+ /// Alias to src.Sem()
+ const sem::Info& sem = src.Sem();
+ /// Name of the uniform buffer variable that holds the array lengths
+ Symbol array_lengths_var;
+ /// A map of pointer-parameter to the name of the new array-length parameter.
+ Hashmap<const sem::Parameter*, Symbol, 8> param_lengths;
+ /// Indices into the uniform buffer array indices that are statically used.
+ std::unordered_set<uint32_t> used_size_indices;
+ /// A vector of array-length parameters which need corresponding array-length arguments for all
+ /// callsites.
+ UniqueVector<const sem::Parameter*, 8> len_params_needing_args;
};
Transform::ApplyResult ArrayLengthFromUniform::Apply(const Program& src,
diff --git a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
index 2a9c6ac..47aedb5 100644
--- a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
@@ -130,17 +130,17 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@compute @workgroup_size(1)
fn main() {
- var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 4u);
+ var len : u32 = (tint_array_lengths.array_lengths[0u][0u] / 4u);
}
)";
@@ -174,11 +174,11 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
struct SB {
x : i32,
@@ -189,7 +189,7 @@
@compute @workgroup_size(1)
fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+ var len : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
}
)";
@@ -225,11 +225,11 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
struct SB {
x : i32,
@@ -240,7 +240,7 @@
@compute @workgroup_size(1)
fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+ var len : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
}
)";
@@ -290,11 +290,11 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 2u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 2u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
struct SB1 {
x : i32,
@@ -323,21 +323,21 @@
@compute @workgroup_size(1)
fn main() {
- var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
- var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u);
- var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
- var len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u);
- var len5 : u32 = (tint_symbol_1.buffer_size[1u][0u] / 16u);
+ var len1 : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
+ var len2 : u32 = ((tint_array_lengths.array_lengths[0u][1u] - 16u) / 16u);
+ var len3 : u32 = (tint_array_lengths.array_lengths[0u][2u] / 16u);
+ var len4 : u32 = ((tint_array_lengths.array_lengths[0u][3u] - 16u) / 16u);
+ var len5 : u32 = (tint_array_lengths.array_lengths[1u][0u] / 16u);
var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
}
)";
ArrayLengthFromUniform::Config cfg({0, 30u});
- cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 2u}, 0);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{1u, 2u}, 1);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{2u, 2u}, 2);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{3u, 2u}, 3);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{4u, 2u}, 4);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 2}, 0);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{1u, 2}, 1);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{2u, 2}, 2);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{3u, 2}, 3);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{4u, 2}, 4);
DataMap data;
data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
@@ -379,11 +379,11 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
struct SB1 {
x : i32,
@@ -412,18 +412,18 @@
@compute @workgroup_size(1)
fn main() {
- var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
- var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
+ var len1 : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
+ var len3 : u32 = (tint_array_lengths.array_lengths[0u][2u] / 16u);
var x : u32 = (len1 + len3);
}
)";
ArrayLengthFromUniform::Config cfg({0, 30u});
- cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 2u}, 0);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{1u, 2u}, 1);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{2u, 2u}, 2);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{3u, 2u}, 3);
- cfg.bindpoint_to_size_index.emplace(BindingPoint{4u, 2u}, 4);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 2}, 0);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{1u, 2}, 1);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{2u, 2}, 2);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{3u, 2}, 3);
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{4u, 2}, 4);
DataMap data;
data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
@@ -486,12 +486,13 @@
}
)";
- auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+ auto* expect =
+ R"(
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
struct SB1 {
x : i32,
@@ -509,7 +510,7 @@
@compute @workgroup_size(1)
fn main() {
- var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+ var len1 : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
var len2 : u32 = arrayLength(&(sb2.arr2));
var x : u32 = (len1 + len2);
}
@@ -544,15 +545,15 @@
)";
auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>,
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
}
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
@compute @workgroup_size(1)
fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+ var len : u32 = ((tint_array_lengths.array_lengths[0u][0u] - 4u) / 4u);
}
@group(0) @binding(0) var<storage, read> sb : SB;
@@ -576,5 +577,133 @@
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
}
+TEST_F(ArrayLengthFromUniformTest, PtrParam_SingleUse) {
+ auto* src = R"(
+@binding(0) @group(0) var<storage, read_write> arr : array<u32>;
+
+fn f2(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return arrayLength(p);
+}
+
+fn f1(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f2(p);
+}
+
+fn f0(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f1(p);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr[0] = f0(&arr);
+}
+)";
+
+ auto* expect =
+ R"(
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 1u>,
+}
+
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
+
+@binding(0) @group(0) var<storage, read_write> arr : array<u32>;
+
+fn f2(p : ptr<storage, array<u32>, read_write>, p_length : u32) -> u32 {
+ return p_length;
+}
+
+fn f1(p : ptr<storage, array<u32>, read_write>, p_length_1 : u32) -> u32 {
+ return f2(p, p_length_1);
+}
+
+fn f0(p : ptr<storage, array<u32>, read_write>, p_length_2 : u32) -> u32 {
+ return f1(p, p_length_2);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr[0] = f0(&(arr), (tint_array_lengths.array_lengths[0u][3u] / 4u));
+}
+)";
+
+ ArrayLengthFromUniform::Config cfg({0, 30u});
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 0}, 3);
+
+ DataMap data;
+ data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(std::unordered_set<uint32_t>({3}),
+ got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
+TEST_F(ArrayLengthFromUniformTest, MissingBindingPoint_PtrParam_MultipleUse) {
+ auto* src = R"(
+@binding(0) @group(0) var<storage, read_write> arr_a : array<u32>;
+@binding(0) @group(1) var<storage, read_write> arr_b : array<u32>;
+
+fn f2(p2 : ptr<storage, array<u32>, read_write>) -> u32 {
+ return arrayLength(p2);
+}
+
+fn f1(p1 : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f2(p1) + arrayLength(p1);
+}
+
+fn f0(p0 : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f1(p0) + arrayLength(p0);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr_a[0] = f0(&arr_a) + arrayLength(&arr_a);
+ arr_b[0] = f0(&arr_b) + arrayLength(&arr_b);
+}
+)";
+
+ auto* expect = R"(
+struct TintArrayLengths {
+ array_lengths : array<vec4<u32>, 2u>,
+}
+
+@group(0) @binding(30) var<uniform> tint_array_lengths : TintArrayLengths;
+
+@binding(0) @group(0) var<storage, read_write> arr_a : array<u32>;
+
+@binding(0) @group(1) var<storage, read_write> arr_b : array<u32>;
+
+fn f2(p2 : ptr<storage, array<u32>, read_write>, p2_length : u32) -> u32 {
+ return p2_length;
+}
+
+fn f1(p1 : ptr<storage, array<u32>, read_write>, p1_length : u32) -> u32 {
+ return (f2(p1, p1_length) + p1_length);
+}
+
+fn f0(p0 : ptr<storage, array<u32>, read_write>, p0_length : u32) -> u32 {
+ return (f1(p0, p0_length) + p0_length);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr_a[0] = (f0(&(arr_a), arrayLength(&(arr_a))) + arrayLength(&(arr_a)));
+ arr_b[0] = (f0(&(arr_b), (tint_array_lengths.array_lengths[1u][1u] / 4u)) + (tint_array_lengths.array_lengths[1u][1u] / 4u));
+}
+)";
+
+ ArrayLengthFromUniform::Config cfg({0, 30u});
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{1, 0}, 5);
+
+ DataMap data;
+ data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::ast::transform
diff --git a/test/tint/bug/chromium/1290107.wgsl.expected.msl b/test/tint/bug/chromium/1290107.wgsl.expected.msl
index de7b6e3..df83392 100644
--- a/test/tint/bug/chromium/1290107.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1290107.wgsl.expected.msl
@@ -14,16 +14,16 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
float f;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const len = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const len = ((*(tint_symbol_1)).array_lengths[0u][0u] / 4u);
return;
}
diff --git a/test/tint/bug/tint/1725.wgsl.expected.msl b/test/tint/bug/tint/1725.wgsl.expected.msl
index a5c54fe..0b157f8 100644
--- a/test/tint/bug/tint/1725.wgsl.expected.msl
+++ b/test/tint/bug/tint/1725.wgsl.expected.msl
@@ -14,23 +14,23 @@
T elements[N];
};
-struct tint_symbol_12 {
+struct tint_symbol_10 {
/* 0x0000 */ tint_array<uint, 1> arr;
};
-struct tint_symbol_7 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
-void tint_symbol_1_inner(uint tint_symbol_2, const device tint_array<uint, 1>* const tint_symbol_9, const constant tint_symbol_7* const tint_symbol_10) {
+void tint_symbol_1_inner(uint tint_symbol_2, const device tint_array<uint, 1>* const tint_symbol_7, const constant TintArrayLengths* const tint_symbol_8) {
int const tint_symbol_3 = 0;
int const tint_symbol_4 = 0;
int const tint_symbol_5 = 0;
- uint const tint_symbol_6 = (*(tint_symbol_9))[min(tint_symbol_2, (((*(tint_symbol_10)).buffer_size[0u][0u] / 4u) - 1u))];
+ uint const tint_symbol_6 = (*(tint_symbol_7))[min(tint_symbol_2, (((*(tint_symbol_8)).array_lengths[0u][0u] / 4u) - 1u))];
}
-kernel void tint_symbol_1(const device tint_symbol_12* tint_symbol_11 [[buffer(0)]], const constant tint_symbol_7* tint_symbol_13 [[buffer(30)]], uint tint_symbol_2 [[thread_index_in_threadgroup]]) {
- tint_symbol_1_inner(tint_symbol_2, &((*(tint_symbol_11)).arr), tint_symbol_13);
+kernel void tint_symbol_1(const device tint_symbol_10* tint_symbol_9 [[buffer(0)]], const constant TintArrayLengths* tint_symbol_11 [[buffer(30)]], uint tint_symbol_2 [[thread_index_in_threadgroup]]) {
+ tint_symbol_1_inner(tint_symbol_2, &((*(tint_symbol_9)).arr), tint_symbol_11);
return;
}
diff --git a/test/tint/bug/tint/2177.wgsl b/test/tint/bug/tint/2177.wgsl
new file mode 100644
index 0000000..eecfa14
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl
@@ -0,0 +1,18 @@
+@binding(0) @group(0) var<storage, read_write> arr : array<u32>;
+
+fn f2(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return arrayLength(p);
+}
+
+fn f1(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f2(p);
+}
+
+fn f0(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f1(p);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr[0] = f0(&arr);
+}
diff --git a/test/tint/bug/tint/2177.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2177.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..79d5ea6
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+RWByteAddressBuffer arr : register(u0);
+
+uint f2_arr() {
+ uint tint_symbol_1 = 0u;
+ arr.GetDimensions(tint_symbol_1);
+ uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ return tint_symbol_2;
+}
+
+uint f1_arr() {
+ return f2_arr();
+}
+
+uint f0_arr() {
+ return f1_arr();
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ arr.Store(0u, asuint(f0_arr()));
+ return;
+}
diff --git a/test/tint/bug/tint/2177.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2177.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..79d5ea6
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+RWByteAddressBuffer arr : register(u0);
+
+uint f2_arr() {
+ uint tint_symbol_1 = 0u;
+ arr.GetDimensions(tint_symbol_1);
+ uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ return tint_symbol_2;
+}
+
+uint f1_arr() {
+ return f2_arr();
+}
+
+uint f0_arr() {
+ return f1_arr();
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ arr.Store(0u, asuint(f0_arr()));
+ return;
+}
diff --git a/test/tint/bug/tint/2177.wgsl.expected.glsl b/test/tint/bug/tint/2177.wgsl.expected.glsl
new file mode 100644
index 0000000..2a6dd66
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer arr_block_ssbo {
+ uint inner[];
+} arr;
+
+uint f2_arr() {
+ return uint(arr.inner.length());
+}
+
+uint f1_arr() {
+ return f2_arr();
+}
+
+uint f0_arr() {
+ return f1_arr();
+}
+
+void tint_symbol() {
+ arr.inner[0] = f0_arr();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/bug/tint/2177.wgsl.expected.msl b/test/tint/bug/tint/2177.wgsl.expected.msl
new file mode 100644
index 0000000..16bd3f5
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.msl
@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct tint_symbol_2 {
+ /* 0x0000 */ tint_array<uint, 1> arr;
+};
+
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
+};
+
+uint f2(device tint_array<uint, 1>* const p, uint p_length) {
+ return p_length;
+}
+
+uint f1(device tint_array<uint, 1>* const p, uint p_length_1) {
+ return f2(p, p_length_1);
+}
+
+uint f0(device tint_array<uint, 1>* const p, uint p_length_2) {
+ return f1(p, p_length_2);
+}
+
+kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(0)]], const constant TintArrayLengths* tint_symbol_3 [[buffer(30)]]) {
+ (*(tint_symbol_1)).arr[0] = f0(&((*(tint_symbol_1)).arr), ((*(tint_symbol_3)).array_lengths[0u][0u] / 4u));
+ return;
+}
+
diff --git a/test/tint/bug/tint/2177.wgsl.expected.spvasm b/test/tint/bug/tint/2177.wgsl.expected.spvasm
new file mode 100644
index 0000000..0e47aa2
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.spvasm
@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %arr_block "arr_block"
+ OpMemberName %arr_block 0 "inner"
+ OpName %arr "arr"
+ OpName %f2_arr "f2_arr"
+ OpName %f1_arr "f1_arr"
+ OpName %f0_arr "f0_arr"
+ OpName %main "main"
+ OpDecorate %arr_block Block
+ OpMemberDecorate %arr_block 0 Offset 0
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpDecorate %arr Binding 0
+ OpDecorate %arr DescriptorSet 0
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %arr_block = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_arr_block = OpTypePointer StorageBuffer %arr_block
+ %arr = OpVariable %_ptr_StorageBuffer_arr_block StorageBuffer
+ %6 = OpTypeFunction %uint
+ %void = OpTypeVoid
+ %16 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %22 = OpConstantNull %int
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %f2_arr = OpFunction %uint None %6
+ %8 = OpLabel
+ %9 = OpArrayLength %uint %arr 0
+ OpReturnValue %9
+ OpFunctionEnd
+ %f1_arr = OpFunction %uint None %6
+ %11 = OpLabel
+ %12 = OpFunctionCall %uint %f2_arr
+ OpReturnValue %12
+ OpFunctionEnd
+ %f0_arr = OpFunction %uint None %6
+ %14 = OpLabel
+ %15 = OpFunctionCall %uint %f1_arr
+ OpReturnValue %15
+ OpFunctionEnd
+ %main = OpFunction %void None %16
+ %19 = OpLabel
+ %24 = OpAccessChain %_ptr_StorageBuffer_uint %arr %uint_0 %22
+ %25 = OpFunctionCall %uint %f0_arr
+ OpStore %24 %25
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2177.wgsl.expected.wgsl b/test/tint/bug/tint/2177.wgsl.expected.wgsl
new file mode 100644
index 0000000..9f03414
--- /dev/null
+++ b/test/tint/bug/tint/2177.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+@binding(0) @group(0) var<storage, read_write> arr : array<u32>;
+
+fn f2(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return arrayLength(p);
+}
+
+fn f1(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f2(p);
+}
+
+fn f0(p : ptr<storage, array<u32>, read_write>) -> u32 {
+ return f1(p);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ arr[0] = f0(&(arr));
+}
diff --git a/test/tint/builtins/arrayLength/complex_via_let.wgsl.expected.msl b/test/tint/builtins/arrayLength/complex_via_let.wgsl.expected.msl
index f92477a..439b1cc 100644
--- a/test/tint/builtins/arrayLength/complex_via_let.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/complex_via_let.wgsl.expected.msl
@@ -14,16 +14,16 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
tint_array<int, 1> a;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.msl b/test/tint/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.msl
index 3fb5295..fb7fba8 100644
--- a/test/tint/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/complex_via_let_no_struct.wgsl.expected.msl
@@ -14,12 +14,12 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_1)).array_lengths[0u][0u] / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/deprecated.wgsl.expected.msl b/test/tint/builtins/arrayLength/deprecated.wgsl.expected.msl
index bc4e7ca..35072c2 100644
--- a/test/tint/builtins/arrayLength/deprecated.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/deprecated.wgsl.expected.msl
@@ -14,17 +14,17 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
tint_array<int, 1> a;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- uint const l2 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ uint const l2 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/simple.wgsl.expected.msl b/test/tint/builtins/arrayLength/simple.wgsl.expected.msl
index f92477a..439b1cc 100644
--- a/test/tint/builtins/arrayLength/simple.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/simple.wgsl.expected.msl
@@ -14,16 +14,16 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
tint_array<int, 1> a;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/simple_no_struct.wgsl.expected.msl b/test/tint/builtins/arrayLength/simple_no_struct.wgsl.expected.msl
index 3fb5295..fb7fba8 100644
--- a/test/tint/builtins/arrayLength/simple_no_struct.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/simple_no_struct.wgsl.expected.msl
@@ -14,12 +14,12 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_1)).array_lengths[0u][0u] / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/via_let.wgsl.expected.msl b/test/tint/builtins/arrayLength/via_let.wgsl.expected.msl
index f92477a..439b1cc 100644
--- a/test/tint/builtins/arrayLength/via_let.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/via_let.wgsl.expected.msl
@@ -14,16 +14,16 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
tint_array<int, 1> a;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/via_let_complex.wgsl.expected.msl b/test/tint/builtins/arrayLength/via_let_complex.wgsl.expected.msl
index f92477a..439b1cc 100644
--- a/test/tint/builtins/arrayLength/via_let_complex.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/via_let_complex.wgsl.expected.msl
@@ -14,16 +14,16 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct S {
tint_array<int, 1> a;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.msl b/test/tint/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.msl
index 3fb5295..fb7fba8 100644
--- a/test/tint/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/via_let_complex_no_struct.wgsl.expected.msl
@@ -14,12 +14,12 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_1)).array_lengths[0u][0u] / 4u);
return;
}
diff --git a/test/tint/builtins/arrayLength/via_let_no_struct.wgsl.expected.msl b/test/tint/builtins/arrayLength/via_let_no_struct.wgsl.expected.msl
index 3fb5295..fb7fba8 100644
--- a/test/tint/builtins/arrayLength/via_let_no_struct.wgsl.expected.msl
+++ b/test/tint/builtins/arrayLength/via_let_no_struct.wgsl.expected.msl
@@ -14,12 +14,12 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
-kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
- uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+kernel void tint_symbol(const constant TintArrayLengths* tint_symbol_1 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_1)).array_lengths[0u][0u] / 4u);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/1588cd.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/1588cd.wgsl.expected.msl
index cc50654..e05a64a 100644
--- a/test/tint/builtins/gen/literal/arrayLength/1588cd.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/1588cd.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<int, 1> arg_0;
};
-void arrayLength_1588cd(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_1588cd(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_1588cd(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_1588cd(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_1588cd(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_1588cd(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_1588cd(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_1588cd(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/61b1c7.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/61b1c7.wgsl.expected.msl
index 21ebe2c..3159a48 100644
--- a/test/tint/builtins/gen/literal/arrayLength/61b1c7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/61b1c7.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<int, 1> arg_0;
};
-void arrayLength_61b1c7(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_61b1c7(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_61b1c7(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_61b1c7(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_61b1c7(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_61b1c7(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_61b1c7(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_61b1c7(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.msl
index d1d5766..4675999 100644
--- a/test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/8421b9.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<half, 1> arg_0;
};
-void arrayLength_8421b9(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 2u);
- *(tint_symbol_4) = res;
+void arrayLength_8421b9(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 2u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_8421b9(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_8421b9(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_8421b9(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_8421b9(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_8421b9(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_8421b9(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/a0f5ca.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/a0f5ca.wgsl.expected.msl
index 41edcde..e4750df 100644
--- a/test/tint/builtins/gen/literal/arrayLength/a0f5ca.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/a0f5ca.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<float, 1> arg_0;
};
-void arrayLength_a0f5ca(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_a0f5ca(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_a0f5ca(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_a0f5ca(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_a0f5ca(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_a0f5ca(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_a0f5ca(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_a0f5ca(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.msl
index 1d5314b..e974459 100644
--- a/test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/cbd6b5.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<half, 1> arg_0;
};
-void arrayLength_cbd6b5(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 2u);
- *(tint_symbol_4) = res;
+void arrayLength_cbd6b5(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 2u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cbd6b5(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cbd6b5(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_cbd6b5(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_cbd6b5(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_cbd6b5(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_cbd6b5(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/cdd123.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/cdd123.wgsl.expected.msl
index 72bd980..3be6c03 100644
--- a/test/tint/builtins/gen/literal/arrayLength/cdd123.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/cdd123.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<float, 1> arg_0;
};
-void arrayLength_cdd123(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_cdd123(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cdd123(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cdd123(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_cdd123(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_cdd123(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_cdd123(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_cdd123(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/cfca0a.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/cfca0a.wgsl.expected.msl
index aa0827f..394279c 100644
--- a/test/tint/builtins/gen/literal/arrayLength/cfca0a.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/cfca0a.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<uint, 1> arg_0;
};
-void arrayLength_cfca0a(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_cfca0a(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cfca0a(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cfca0a(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_cfca0a(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_cfca0a(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_cfca0a(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_cfca0a(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/literal/arrayLength/eb510f.wgsl.expected.msl b/test/tint/builtins/gen/literal/arrayLength/eb510f.wgsl.expected.msl
index 6838fcb..f4a0aa7 100644
--- a/test/tint/builtins/gen/literal/arrayLength/eb510f.wgsl.expected.msl
+++ b/test/tint/builtins/gen/literal/arrayLength/eb510f.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<uint, 1> arg_0;
};
-void arrayLength_eb510f(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_eb510f(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_eb510f(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_eb510f(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_eb510f(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_eb510f(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_eb510f(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_eb510f(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/1588cd.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/1588cd.wgsl.expected.msl
index cc50654..e05a64a 100644
--- a/test/tint/builtins/gen/var/arrayLength/1588cd.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/1588cd.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<int, 1> arg_0;
};
-void arrayLength_1588cd(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_1588cd(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_1588cd(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_1588cd(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_1588cd(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_1588cd(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_1588cd(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_1588cd(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/61b1c7.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/61b1c7.wgsl.expected.msl
index 21ebe2c..3159a48 100644
--- a/test/tint/builtins/gen/var/arrayLength/61b1c7.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/61b1c7.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<int, 1> arg_0;
};
-void arrayLength_61b1c7(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_61b1c7(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_61b1c7(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_61b1c7(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_61b1c7(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_61b1c7(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_61b1c7(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_61b1c7(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.msl
index d1d5766..4675999 100644
--- a/test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/8421b9.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<half, 1> arg_0;
};
-void arrayLength_8421b9(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 2u);
- *(tint_symbol_4) = res;
+void arrayLength_8421b9(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 2u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_8421b9(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_8421b9(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_8421b9(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_8421b9(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_8421b9(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_8421b9(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/a0f5ca.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/a0f5ca.wgsl.expected.msl
index 41edcde..e4750df 100644
--- a/test/tint/builtins/gen/var/arrayLength/a0f5ca.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/a0f5ca.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<float, 1> arg_0;
};
-void arrayLength_a0f5ca(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_a0f5ca(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_a0f5ca(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_a0f5ca(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_a0f5ca(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_a0f5ca(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_a0f5ca(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_a0f5ca(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.msl
index 1d5314b..e974459 100644
--- a/test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/cbd6b5.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<half, 1> arg_0;
};
-void arrayLength_cbd6b5(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 2u);
- *(tint_symbol_4) = res;
+void arrayLength_cbd6b5(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 2u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cbd6b5(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cbd6b5(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_cbd6b5(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_cbd6b5(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_cbd6b5(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_cbd6b5(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/cdd123.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/cdd123.wgsl.expected.msl
index 72bd980..3be6c03 100644
--- a/test/tint/builtins/gen/var/arrayLength/cdd123.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/cdd123.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<float, 1> arg_0;
};
-void arrayLength_cdd123(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_cdd123(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cdd123(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cdd123(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_cdd123(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_cdd123(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_cdd123(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_cdd123(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/cfca0a.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/cfca0a.wgsl.expected.msl
index aa0827f..394279c 100644
--- a/test/tint/builtins/gen/var/arrayLength/cfca0a.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/cfca0a.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RO {
tint_array<uint, 1> arg_0;
};
-void arrayLength_cfca0a(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_cfca0a(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_cfca0a(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_cfca0a(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(0)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
- arrayLength_cfca0a(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(0)]]) {
+ arrayLength_cfca0a(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(0)]]) {
- arrayLength_cfca0a(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(0)]]) {
+ arrayLength_cfca0a(tint_symbol_9, tint_symbol_10);
return;
}
diff --git a/test/tint/builtins/gen/var/arrayLength/eb510f.wgsl.expected.msl b/test/tint/builtins/gen/var/arrayLength/eb510f.wgsl.expected.msl
index 6838fcb..f4a0aa7 100644
--- a/test/tint/builtins/gen/var/arrayLength/eb510f.wgsl.expected.msl
+++ b/test/tint/builtins/gen/var/arrayLength/eb510f.wgsl.expected.msl
@@ -14,42 +14,42 @@
T elements[N];
};
-struct tint_symbol_1 {
- /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+struct TintArrayLengths {
+ /* 0x0000 */ tint_array<uint4, 1> array_lengths;
};
struct SB_RW {
tint_array<uint, 1> arg_0;
};
-void arrayLength_eb510f(const constant tint_symbol_1* const tint_symbol_3, device uint* const tint_symbol_4) {
- uint res = (((*(tint_symbol_3)).buffer_size[0u][0u] - 0u) / 4u);
- *(tint_symbol_4) = res;
+void arrayLength_eb510f(const constant TintArrayLengths* const tint_symbol_1, device uint* const tint_symbol_2) {
+ uint res = (((*(tint_symbol_1)).array_lengths[0u][0u] - 0u) / 4u);
+ *(tint_symbol_2) = res;
}
struct tint_symbol {
float4 value [[position]];
};
-float4 vertex_main_inner(const constant tint_symbol_1* const tint_symbol_5, device uint* const tint_symbol_6) {
- arrayLength_eb510f(tint_symbol_5, tint_symbol_6);
+float4 vertex_main_inner(const constant TintArrayLengths* const tint_symbol_3, device uint* const tint_symbol_4) {
+ arrayLength_eb510f(tint_symbol_3, tint_symbol_4);
return float4(0.0f);
}
-vertex tint_symbol vertex_main(const constant tint_symbol_1* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
- float4 const inner_result = vertex_main_inner(tint_symbol_7, tint_symbol_8);
+vertex tint_symbol vertex_main(const constant TintArrayLengths* tint_symbol_5 [[buffer(30)]], device uint* tint_symbol_6 [[buffer(1)]]) {
+ float4 const inner_result = vertex_main_inner(tint_symbol_5, tint_symbol_6);
tint_symbol wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}
-fragment void fragment_main(const constant tint_symbol_1* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
- arrayLength_eb510f(tint_symbol_9, tint_symbol_10);
+fragment void fragment_main(const constant TintArrayLengths* tint_symbol_7 [[buffer(30)]], device uint* tint_symbol_8 [[buffer(1)]]) {
+ arrayLength_eb510f(tint_symbol_7, tint_symbol_8);
return;
}
-kernel void compute_main(const constant tint_symbol_1* tint_symbol_11 [[buffer(30)]], device uint* tint_symbol_12 [[buffer(1)]]) {
- arrayLength_eb510f(tint_symbol_11, tint_symbol_12);
+kernel void compute_main(const constant TintArrayLengths* tint_symbol_9 [[buffer(30)]], device uint* tint_symbol_10 [[buffer(1)]]) {
+ arrayLength_eb510f(tint_symbol_9, tint_symbol_10);
return;
}
