Import Tint changes from Dawn
Changes:
- c8fb20b38ba33dcaf4a7d688c5508a4b729033d0 [tint][msl] Fix doxygen warning / error by Ben Clayton <bclayton@google.com>
- f0057a56feedf3c617c9544a65416460c4f4aa9d Add inspector helper for texture information. by dan sinclair <dsinclair@chromium.org>
- ffd485c685040edb1e678165dcbf0e841cfa0298 [tint][msl] Replace volatile loop conditional with outer ... by Ben Clayton <bclayton@google.com>
- 267845cb1db12737854d8b12e900711d10830b45 [spirv-reader] Parse IO struct member decorations by James Price <jrprice@google.com>
- a3cad6827d5ce6d6d437ae36dd5dce7924d5dee9 [spirv-reader] Parse Invariant decoration on var by James Price <jrprice@google.com>
- 3d63299556a1899ef84df2d7d811caf02a855349 [spirv-reader] Parse user-defined IO decorations by James Price <jrprice@google.com>
- fe068476847cba50af47c04d50065bc9ac22c3ec [spirv-reader] Parse input/output builtins by James Price <jrprice@google.com>
- 9be037cf1bcbd4a37d083d079f4e457e2b7f5182 [tint][msl]: Use a volatile bool to control loop preserva... by Ben Clayton <bclayton@google.com>
- 9cfc80a7865ae17e71631cfe481a8d067a3dcc85 [spirv-reader] Handle StorageBuffer storage class by James Price <jrprice@google.com>
- ea0fda3e68e565d4f4610b4521f6cafe663bfa84 [ast,msl]: packed_vec3: Unroll loop for small composite c... by David Neto <dneto@google.com>
GitOrigin-RevId: c8fb20b38ba33dcaf4a7d688c5508a4b729033d0
Change-Id: Ia9dd2769bf321ad75f64b6397bd8f163c7c67ccf
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/170521
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/type/manager.h b/src/tint/lang/core/type/manager.h
index 0e8b324..6c12a6c 100644
--- a/src/tint/lang/core/type/manager.h
+++ b/src/tint/lang/core/type/manager.h
@@ -65,13 +65,13 @@
/// @returns the default access control for a memory view with the given address space.
static constexpr inline core::Access DefaultAccessFor(core::AddressSpace space) {
switch (space) {
+ case core::AddressSpace::kIn:
case core::AddressSpace::kPushConstant:
case core::AddressSpace::kUniform:
case core::AddressSpace::kHandle:
return core::Access::kRead;
case core::AddressSpace::kUndefined:
- case core::AddressSpace::kIn:
case core::AddressSpace::kOut:
case core::AddressSpace::kFunction:
case core::AddressSpace::kPixelLocal:
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 63dcf81..718770e 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -2122,8 +2122,7 @@
};
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
- Line() << "while (true) {";
- EmitLoopPreserver();
+ Line() << IsolateUB() << " while(true) {";
{
ScopedIndent si(this);
if (!EmitStatements(stmt->body->statements)) {
@@ -2193,8 +2192,7 @@
};
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
- Line() << "while (true) {";
- EmitLoopPreserver();
+ Line() << IsolateUB() << " while(true) {";
IncrementIndent();
TINT_DEFER({
DecrementIndent();
@@ -2217,7 +2215,7 @@
// For-loop can be generated.
{
auto out = Line();
- out << "for";
+ out << IsolateUB() << " for";
{
ScopedParen sp(out);
@@ -2235,7 +2233,6 @@
}
out << " {";
}
- EmitLoopPreserver();
{
auto emit_continuing = [] { return true; };
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
@@ -2268,8 +2265,7 @@
// as a regular while in MSL. Instead we need to generate a `while(true)` loop.
bool emit_as_loop = cond_pre.lines.size() > 0;
if (emit_as_loop) {
- Line() << "while (true) {";
- EmitLoopPreserver();
+ Line() << IsolateUB() << " while(true) {";
IncrementIndent();
TINT_DEFER({
DecrementIndent();
@@ -2283,16 +2279,7 @@
}
} else {
// While can be generated.
- {
- auto out = Line();
- out << "while";
- {
- ScopedParen sp(out);
- out << cond_buf.str();
- }
- out << " {";
- }
- EmitLoopPreserver();
+ Line() << IsolateUB() << " while(" << cond_buf.str() << ") {";
if (!EmitStatementsWithIndent(stmt->body->statements)) {
return false;
}
@@ -3031,32 +3018,16 @@
return true;
}
-void ASTPrinter::EmitLoopPreserver() {
- IncrementIndent();
- // This statement prevents the MSL compiler from erasing a loop during
- // optimizations. In the AIR dialiect of LLVM IR, WGSL loops should compile
- // to a loop that contains an 'asm' call with a 'sideeffect' annotation.
- //
- // For example, compile a WGSL file with a trivial while(1) loop to 'a.metal',
- // then compile that to AIR (LLVM IR dialect):
- //
- // xcrun metal a.metal -S -o -
- //
- // The loop in the AIR should look something like this:
- //
- // 1: ...
- // br label %2
- //
- // 2: ; preds = %1, %2
- // tail call void asm sideeffect "", ""() #1, !srcloc !27
- // br label %2, !llvm.loop !28
- //
- // It is important that the 'sideeffect' annotation exist. That tells the
- // optimizer that the instruction has side effects invisible to the
- // optimizer, and therefore the loop should not be eliminated.
- Line() << R"(__asm__("");)";
-
- DecrementIndent();
+std::string_view ASTPrinter::IsolateUB() {
+ if (isolate_ub_macro_name_.empty()) {
+ isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
+ auto volatile_true = UniqueIdentifier("tint_volatile_true");
+ Line(&helpers_) << "#define " << isolate_ub_macro_name_ << " \\";
+ Line(&helpers_) << " if (volatile bool " << volatile_true << " = true; " << volatile_true
+ << ")";
+ Line(&helpers_);
+ }
+ return isolate_ub_macro_name_;
}
template <typename F>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index af251db..70269c9 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -378,9 +378,10 @@
const ast::CallExpression* expr,
const sem::BuiltinFn* builtin);
- /// Emits a code sequence that preserves a loop during
- /// optimizations even if the loop is infinite.
- void EmitLoopPreserver();
+ /// Lazilly generates the TINT_ISOLATE_UB macro, used to prevent UB statements from affecting
+ /// later logic.
+ /// @return the unique name of the TINT_ISOLATE_UB macro.
+ std::string_view IsolateUB();
/// Handles generating a builtin name
/// @param builtin the semantic info for the builtin
@@ -437,6 +438,10 @@
std::function<bool()> emit_continuing_;
+ /// The name of the macro used to prevent UB affecting later control flow.
+ /// Do not use this directly, instead call IsolateUB().
+ std::string isolate_ub_macro_name_;
+
/// Name of atomicCompareExchangeWeak() helper for the given pointer storage
/// class and struct return type
using ACEWKeyType =
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index 0477e48..74d8fdd 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -226,13 +226,15 @@
T elements[N];
};
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
struct tint_symbol_3 {
tint_array<float2x2, 4> m;
};
void comp_main_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
- for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
- __asm__("");
+ TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
uint const i = idx;
(*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
}
diff --git a/src/tint/lang/msl/writer/ast_printer/continue_test.cc b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
index a79ddc3..727f05c 100644
--- a/src/tint/lang/msl/writer/ast_printer/continue_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
@@ -39,16 +39,24 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
+ ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+ EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
- ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+kernel void test_function() {
+ TINT_ISOLATE_UB while(true) {
if (false) {
break;
}
continue;
}
+ return;
+}
+
)");
}
diff --git a/src/tint/lang/msl/writer/ast_printer/loop_test.cc b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
index 450592b..7aed227 100644
--- a/src/tint/lang/msl/writer/ast_printer/loop_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
@@ -46,13 +46,21 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
+ ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+ EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
- ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+fragment void F() {
+ TINT_ISOLATE_UB while(true) {
break;
}
+ return;
+}
+
)");
}
@@ -67,16 +75,27 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
+ ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+ EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
- ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+void a_statement() {
+}
+
+fragment void F() {
+ TINT_ISOLATE_UB while(true) {
break;
{
a_statement();
}
}
+ return;
+}
+
)");
}
@@ -91,17 +110,28 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
+ ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+ EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
- ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+ if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+void a_statement() {
+}
+
+fragment void F() {
+ TINT_ISOLATE_UB while(true) {
break;
{
a_statement();
if (true) { break; }
}
}
+ return;
+}
+
)");
}
@@ -125,23 +155,19 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
- while (true) {
- __asm__("");
- break;
- {
- a_statement();
- }
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ TINT_ISOLATE_UB while(true) {
+ break;
{
- lhs = rhs;
- if (true) { break; }
+ a_statement();
}
}
+ {
+ lhs = rhs;
+ if (true) { break; }
+ }
+}
)");
}
@@ -167,18 +193,15 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
- float lhs = 2.5f;
- float other = 0.0f;
- break;
- {
- lhs = rhs;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ float lhs = 2.5f;
+ float other = 0.0f;
+ break;
+ {
+ lhs = rhs;
}
+}
)");
}
@@ -193,13 +216,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( for(; ; ) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; ; ) {
+ return;
+}
)");
}
@@ -214,13 +234,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( for(int i = 0; ; ) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(int i = 0; ; ) {
+ return;
+}
)");
}
@@ -243,19 +260,16 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( {
- {
- f(1);
- f(2);
- }
- for(; ; ) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"({
+ {
+ f(1);
+ f(2);
}
+ TINT_ISOLATE_UB for(; ; ) {
+ return;
+ }
+}
)");
}
@@ -270,13 +284,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( for(; true; ) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; true; ) {
+ return;
+}
)");
}
@@ -292,14 +303,11 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
EXPECT_EQ(gen.Result(),
- R"( for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
- __asm__("");
- return;
- }
+ R"(TINT_ISOLATE_UB for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ return;
+}
)");
}
@@ -322,17 +330,14 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while (true) {
- __asm__("");
- return;
- {
- f(1);
- f(2);
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ return;
+ {
+ f(1);
+ f(2);
}
+}
)");
}
@@ -349,14 +354,12 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(),
- R"( for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
- __asm__("");
- a_statement();
- }
+ EXPECT_EQ(
+ gen.Result(),
+ R"(TINT_ISOLATE_UB for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ a_statement();
+}
)");
}
@@ -380,24 +383,21 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( {
+ EXPECT_EQ(gen.Result(), R"({
+ {
+ f(1);
+ f(2);
+ }
+ TINT_ISOLATE_UB while(true) {
+ if (!(true)) { break; }
+ return;
{
- f(1);
- f(2);
- }
- while (true) {
- __asm__("");
- if (!(true)) { break; }
- return;
- {
- f(3);
- f(4);
- }
+ f(3);
+ f(4);
}
}
+}
)");
}
@@ -411,13 +411,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while(true) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ return;
+}
)");
}
@@ -431,13 +428,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while(true) {
- __asm__("");
- continue;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+ continue;
+}
)");
}
@@ -454,13 +448,10 @@
ASTPrinter& gen = Build();
- gen.IncrementIndent();
-
ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
- EXPECT_EQ(gen.Result(), R"( while((t && false)) {
- __asm__("");
- return;
- }
+ EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while((t && false)) {
+ return;
+}
)");
}
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
index 9d651c6..a61e3f8 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
@@ -58,6 +58,10 @@
namespace tint::msl::writer {
+/// Arrays larger than this will be packed/unpacked with a for loop.
+/// Arrays up to this size will be packed/unpacked with a sequence of statements.
+static constexpr uint32_t kMaxSeriallyUnpackedArraySize = 8;
+
/// PIMPL state for the transform
struct PackedVec3::State {
/// Constructor
@@ -242,7 +246,6 @@
const std::function<ast::Type()>& out_type) {
// Allocate a variable to hold the return value of the function.
tint::Vector<const ast::Statement*, 4> statements;
- statements.Push(b.Decl(b.Var("result", out_type())));
// Helper that generates a loop to copy and pack/unpack elements of an array to the result:
// for (var i = 0u; i < num_elements; i = i + 1) {
@@ -250,15 +253,28 @@
// }
auto copy_array_elements = [&](uint32_t num_elements,
const core::type::Type* element_type) {
- // Generate an expression for packing or unpacking an element of the array.
- auto* element = pack_or_unpack_element(b.IndexAccessor("in", "i"), element_type);
- statements.Push(b.For( //
- b.Decl(b.Var("i", b.ty.u32())), //
- b.LessThan("i", u32(num_elements)), //
- b.Assign("i", b.Add("i", 1_a)), //
- b.Block(tint::Vector{
- b.Assign(b.IndexAccessor("result", "i"), element),
- })));
+ // Generate code for unpacking the array.
+ if (num_elements <= kMaxSeriallyUnpackedArraySize) {
+ // Generate a variable with an explicit initializer.
+ tint::Vector<const ast::Expression*, 8> elements;
+ for (uint32_t i = 0; i < num_elements; i++) {
+ elements.Push(pack_or_unpack_element(
+ b.IndexAccessor("in", b.Expr(core::AInt(i))), element_type));
+ }
+ statements.Push(b.Decl(b.Var("result", b.Call(out_type(), b.ExprList(elements)))));
+ } else {
+ statements.Push(b.Decl(b.Var("result", out_type())));
+ // Generate a for loop.
+ // Generate an expression for packing or unpacking an element of the array.
+ auto* element = pack_or_unpack_element(b.IndexAccessor("in", "i"), element_type);
+ statements.Push(b.For( //
+ b.Decl(b.Var("i", b.ty.u32())), //
+ b.LessThan("i", u32(num_elements)), //
+ b.Assign("i", b.Add("i", 1_a)), //
+ b.Block(tint::Vector{
+ b.Assign(b.IndexAccessor("result", "i"), element),
+ })));
+ }
};
// Copy the elements of the value over to the result.
@@ -272,6 +288,7 @@
copy_array_elements(mat->columns(), mat->ColumnType());
},
[&](const core::type::Struct* str) {
+ statements.Push(b.Decl(b.Var("result", out_type())));
// Copy the struct members over one at a time, packing/unpacking as necessary.
for (auto* member : str->Members()) {
const ast::Expression* element =
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
index 220dabb..cc002ba 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
@@ -502,7 +502,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray) {
+TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray_Small) {
auto* src = R"(
@group(0) @binding(0) var<storage> arr : array<vec3<f32>, 4>;
@@ -520,14 +520,49 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
+ return result;
+}
+
+@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 4u>;
+
+fn f() {
+ let x = tint_unpack_vec3_in_composite(arr);
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray_Large) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage> arr : array<vec3<f32>, 9>;
+
+fn f() {
+ let x = arr;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
result[i] = vec3<f32>(in[i].elements);
}
return result;
}
-@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 4u>;
+@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 9u>;
fn f() {
let x = tint_unpack_vec3_in_composite(arr);
@@ -630,7 +665,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS_Small) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 2>;
@@ -648,10 +683,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
- var result : array<tint_packed_vec3_f32_array_element, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
return result;
}
@@ -668,6 +700,52 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS_Large) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 9>;
+
+fn f() {
+ arr = array(vec3(1.5, 2.5, 3.5),
+ vec3(4.5, 5.5, 6.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5),
+ vec3(7.5, 8.5, 9.5));
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+ var result : array<tint_packed_vec3_f32_array_element, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+ }
+ return result;
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+fn f() {
+ arr = tint_pack_vec3_in_composite(array(vec3(1.5, 2.5, 3.5), vec3(4.5, 5.5, 6.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5)));
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_RefRHS) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 2>;
@@ -731,7 +809,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS) {
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS_Small) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 4>;
@group(0) @binding(1) var<uniform> in_arr : array<vec3<f32>, 4>;
@@ -769,6 +847,44 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS_Large) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 9>;
+@group(0) @binding(1) var<uniform> in_arr : array<vec3<f32>, 9>;
+@group(0) @binding(2) var<uniform> in_vec : vec3<f32>;
+
+fn f() {
+ arr[0] = in_arr[0];
+ arr[1] = in_vec;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(1) var<uniform> in_arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(2) var<uniform> in_vec : __packed_vec3<f32>;
+
+fn f() {
+ arr[0].elements = in_arr[0].elements;
+ arr[1].elements = in_vec;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, ArrayOfVec3_WriteComponent_MemberAccessor) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 4>;
@@ -847,10 +963,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -975,10 +1088,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -1156,7 +1266,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray) {
+TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray_Small) {
auto* src = R"(
@group(0) @binding(0) var<storage> arr : array<mat3x3<f32>, 4>;
@@ -1174,22 +1284,59 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
+ return result;
+}
+
+@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
+
+fn f() {
+ let x = tint_unpack_vec3_in_composite_1(arr);
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray_Large) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage> arr : array<mat3x3<f32>, 9>;
+
+fn f() {
+ let x = arr;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+ var result : array<mat3x3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
result[i] = tint_unpack_vec3_in_composite(in[i]);
}
return result;
}
-@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
+@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
fn f() {
let x = tint_unpack_vec3_in_composite_1(arr);
@@ -1220,10 +1367,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -1330,7 +1474,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS_Small) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 2>;
@@ -1348,18 +1492,12 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
return result;
}
@@ -1376,6 +1514,57 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS_Large) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 9>;
+
+fn f() {
+ arr = array(mat3x3<f32>(),
+ mat3x3(1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5, 8.5, 9.5),
+ mat3x3f(),
+ mat3x3f(),
+ mat3x3f(),
+ mat3x3f(),
+ mat3x3f(),
+ mat3x3f(),
+ mat3x3f());
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
+ return result;
+}
+
+fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 9u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 9u> {
+ var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_pack_vec3_in_composite(in[i]);
+ }
+ return result;
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+
+fn f() {
+ arr = tint_pack_vec3_in_composite_1(array(mat3x3<f32>(), mat3x3(1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5, 8.5, 9.5), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f()));
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_RefRHS) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 2>;
@@ -1427,10 +1616,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -2025,7 +2211,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct_Small) {
auto* src = R"(
struct S {
arr : array<vec3<f32>, 4>,
@@ -2052,10 +2238,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
@@ -2082,7 +2265,64 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct_Large) {
+ auto* src = R"(
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = vec3<f32>(in[i].elements);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : S_tint_packed_vec3) -> S {
+ var result : S;
+ result.arr = tint_unpack_vec3_in_composite(in.arr);
+ return result;
+}
+
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+ let x = tint_unpack_vec3_in_composite_1(P);
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray_Small) {
auto* src = R"(
struct S {
arr : array<vec3<f32>, 4>,
@@ -2109,15 +2349,63 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
+ return result;
+}
+
+struct S {
+ arr : array<vec3<f32>, 4>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+ let x = tint_unpack_vec3_in_composite(P.arr);
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray_Large) {
+ auto* src = R"(
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P.arr;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
result[i] = vec3<f32>(in[i].elements);
}
return result;
}
struct S {
- arr : array<vec3<f32>, 4>,
+ arr : array<vec3<f32>, 9>,
}
@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
@@ -2262,7 +2550,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS_Small) {
auto* src = R"(
struct S {
arr : array<vec3<f32>, 2>,
@@ -2290,10 +2578,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
- var result : array<tint_packed_vec3_f32_array_element, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
return result;
}
@@ -2320,6 +2605,72 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS_Large) {
+ auto* src = R"(
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P = S(array(vec3(1.5, 4.5, 7.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5)));
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+ var result : array<tint_packed_vec3_f32_array_element, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+ }
+ return result;
+}
+
+fn tint_pack_vec3_in_composite_1(in : S) -> S_tint_packed_vec3 {
+ var result : S_tint_packed_vec3;
+ result.arr = tint_pack_vec3_in_composite(in.arr);
+ return result;
+}
+
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S_tint_packed_vec3;
+
+fn f() {
+ P = tint_pack_vec3_in_composite_1(S(array(vec3(1.5, 4.5, 7.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5))));
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_RefRHS) {
auto* src = R"(
struct S {
@@ -2366,7 +2717,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS_Small) {
auto* src = R"(
struct S {
arr : array<vec3<f32>, 2>,
@@ -2393,10 +2744,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
- var result : array<tint_packed_vec3_f32_array_element, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
return result;
}
@@ -2417,6 +2765,65 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS_Large) {
+ auto* src = R"(
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+ P.arr = array(vec3(1.5, 4.5, 7.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5),
+ vec3(9.5, 6.5, 3.5));
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+ var result : array<tint_packed_vec3_f32_array_element, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+ }
+ return result;
+}
+
+struct S {
+ arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S_tint_packed_vec3;
+
+fn f() {
+ P.arr = tint_pack_vec3_in_composite(array(vec3(1.5, 4.5, 7.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5)));
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_RefRHS) {
auto* src = R"(
struct S {
@@ -2684,10 +3091,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -2741,10 +3145,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -2922,10 +3323,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -3025,10 +3423,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -3289,7 +3684,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct_Small) {
auto* src = R"(
struct S {
arr : array<mat3x3<f32>, 4>,
@@ -3316,18 +3711,12 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
@@ -3354,10 +3743,72 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct_Large) {
+ auto* src = R"(
+struct S {
+ arr : array<mat3x3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+ let x = P;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+ var result : array<mat3x3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : S_tint_packed_vec3) -> S {
+ var result : S;
+ result.arr = tint_unpack_vec3_in_composite_1(in.arr);
+ return result;
+}
+
+struct S {
+ arr : array<mat3x3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+ let x = tint_unpack_vec3_in_composite_2(P);
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadArray) {
auto* src = R"(
struct S {
- arr : array<mat3x3<f32>, 4>,
+ arr : array<mat3x3<f32>, 9>,
}
@group(0) @binding(0) var<storage> P : S;
@@ -3377,27 +3828,24 @@
struct S_tint_packed_vec3 {
@align(16)
- arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>,
+ arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
-fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+ var result : array<mat3x3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
result[i] = tint_unpack_vec3_in_composite(in[i]);
}
return result;
}
struct S {
- arr : array<mat3x3<f32>, 4>,
+ arr : array<mat3x3<f32>, 9>,
}
@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
@@ -3440,10 +3888,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -3621,18 +4066,12 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
return result;
}
@@ -3732,18 +4171,12 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
- for(var i : u32; (i < 2u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
return result;
}
@@ -3846,10 +4279,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -4510,10 +4940,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
@@ -4528,10 +4955,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -4759,10 +5183,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
@@ -4965,18 +5386,12 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -5047,18 +5462,12 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -5120,34 +5529,22 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -5213,26 +5610,17 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
@@ -5246,26 +5634,17 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
fn tint_pack_vec3_in_composite_2(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -5460,18 +5839,12 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -5551,18 +5924,12 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -5635,34 +6002,22 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -5741,26 +6096,17 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
@@ -5774,26 +6120,17 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
fn tint_pack_vec3_in_composite_2(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -5868,10 +6205,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -5947,10 +6281,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -6041,10 +6372,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -6133,26 +6461,17 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
@@ -6239,26 +6558,17 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
@@ -6484,10 +6794,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -6706,10 +7013,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -6933,10 +7237,7 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
@@ -7000,18 +7301,12 @@
}
fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
- var result : array<tint_packed_vec3_f32_array_element, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<array<vec3<f32>, 4u>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 4u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -7075,10 +7370,7 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
@@ -7140,18 +7432,12 @@
}
fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
- var result : array<tint_packed_vec3_f32_array_element, 3u>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
- }
+ var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
return result;
}
fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
- var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_pack_vec3_in_composite(in[i]);
- }
+ var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
return result;
}
@@ -7222,42 +7508,27 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
- var result : array<array<vec3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_2(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
- var result : array<array<mat3x3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_3(in[i]);
- }
+ var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
return result;
}
@@ -7269,10 +7540,7 @@
}
fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
- var result : array<S, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_5(in[i]);
- }
+ var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
return result;
}
@@ -7304,7 +7572,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, NestedArrays_VarInit) {
+TEST_F(PackedVec3Test, NestedArrays_VarInit_Small) {
auto* src = R"(
struct S {
arr_v : array<array<vec3<f32>, 4>, 4>,
@@ -7345,42 +7613,27 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
- var result : array<array<vec3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_2(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
- var result : array<array<mat3x3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_3(in[i]);
- }
+ var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
return result;
}
@@ -7392,10 +7645,7 @@
}
fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
- var result : array<S, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_5(in[i]);
- }
+ var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
return result;
}
@@ -7427,7 +7677,127 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, NestedArrays_VarAssignment) {
+TEST_F(PackedVec3Test, NestedArrays_VarInit_Large) {
+ auto* src = R"(
+struct S {
+ arr_v : array<array<vec3<f32>, 9>, 9>,
+ arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S, 9>;
+
+fn f() {
+ var full_var : array<S, 9> = arr_s;
+ var struct_var : S = arr_s[0];
+ var outer_arr_v_var : array<array<vec3<f32>, 9>, 9> = arr_s[0].arr_v;
+ var inner_arr_v_var : array<vec3<f32>, 9> = arr_s[0].arr_v[1];
+ var v_var : vec3<f32> = arr_s[0].arr_v[1][2];
+ var v_element_var : f32 = arr_s[0].arr_v[1][2].y;
+ var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9> = arr_s[0].arr_m;
+ var inner_arr_m_var : array<mat3x3<f32>, 9> = arr_s[0].arr_m[1];
+ var m_var : mat3x3<f32> = arr_s[0].arr_m[1][2];
+ var m_col_var : vec3<f32> = arr_s[0].arr_m[1][2][0];
+ var m_element_var : f32 = arr_s[0].arr_m[1][2][0].y;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr_v : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+ @align(16)
+ arr_m : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = vec3<f32>(in[i].elements);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>) -> array<array<vec3<f32>, 9u>, 9u> {
+ var result : array<array<vec3<f32>, 9u>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+ var result : array<mat3x3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_2(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>) -> array<array<mat3x3<f32>, 9u>, 9u> {
+ var result : array<array<mat3x3<f32>, 9u>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_3(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_5(in : S_tint_packed_vec3) -> S {
+ var result : S;
+ result.arr_v = tint_unpack_vec3_in_composite_1(in.arr_v);
+ result.arr_m = tint_unpack_vec3_in_composite_4(in.arr_m);
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 9u>) -> array<S, 9u> {
+ var result : array<S, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_5(in[i]);
+ }
+ return result;
+}
+
+struct S {
+ arr_v : array<array<vec3<f32>, 9>, 9>,
+ arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S_tint_packed_vec3, 9u>;
+
+fn f() {
+ var full_var : array<S, 9> = tint_unpack_vec3_in_composite_6(arr_s);
+ var struct_var : S = tint_unpack_vec3_in_composite_5(arr_s[0]);
+ var outer_arr_v_var : array<array<vec3<f32>, 9>, 9> = tint_unpack_vec3_in_composite_1(arr_s[0].arr_v);
+ var inner_arr_v_var : array<vec3<f32>, 9> = tint_unpack_vec3_in_composite(arr_s[0].arr_v[1]);
+ var v_var : vec3<f32> = vec3<f32>(arr_s[0].arr_v[1][2].elements);
+ var v_element_var : f32 = arr_s[0].arr_v[1][2].elements.y;
+ var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9> = tint_unpack_vec3_in_composite_4(arr_s[0].arr_m);
+ var inner_arr_m_var : array<mat3x3<f32>, 9> = tint_unpack_vec3_in_composite_3(arr_s[0].arr_m[1]);
+ var m_var : mat3x3<f32> = tint_unpack_vec3_in_composite_2(arr_s[0].arr_m[1][2]);
+ var m_col_var : vec3<f32> = vec3<f32>(arr_s[0].arr_m[1][2][0].elements);
+ var m_element_var : f32 = arr_s[0].arr_m[1][2][0].elements.y;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, NestedArrays_VarAssignment_Small) {
auto* src = R"(
struct S {
arr_v : array<array<vec3<f32>, 4>, 4>,
@@ -7480,42 +7850,27 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
- var result : array<array<vec3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite(in[i]);
- }
+ var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
- var result : array<mat3x3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_2(in[i]);
- }
+ var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
return result;
}
fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
- var result : array<array<mat3x3<f32>, 4u>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_3(in[i]);
- }
+ var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
return result;
}
@@ -7527,10 +7882,7 @@
}
fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
- var result : array<S, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = tint_unpack_vec3_in_composite_5(in[i]);
- }
+ var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
return result;
}
@@ -7573,6 +7925,149 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, NestedArrays_VarAssignment_Large) {
+ auto* src = R"(
+struct S {
+ arr_v : array<array<vec3<f32>, 9>, 9>,
+ arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S, 9>;
+
+fn f() {
+ var full_var : array<S, 9>;
+ var struct_var : S;
+ var outer_arr_v_var : array<array<vec3<f32>, 9>, 9>;
+ var inner_arr_v_var : array<vec3<f32>, 9>;
+ var v_var : vec3<f32>;
+ var v_element_var : f32;
+ var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9>;
+ var inner_arr_m_var : array<mat3x3<f32>, 9>;
+ var m_var : mat3x3<f32>;
+ var m_col_var : vec3<f32>;
+ var m_element_var : f32;
+
+ full_var = arr_s;
+ struct_var = arr_s[0];
+ outer_arr_v_var = arr_s[0].arr_v;
+ inner_arr_v_var = arr_s[0].arr_v[1];
+ v_var = arr_s[0].arr_v[1][2];
+ v_element_var = arr_s[0].arr_v[1][2].y;
+ outer_arr_m_var = arr_s[0].arr_m;
+ inner_arr_m_var = arr_s[0].arr_m[1];
+ m_var = arr_s[0].arr_m[1][2];
+ m_col_var = arr_s[0].arr_m[1][2][0];
+ m_element_var = arr_s[0].arr_m[1][2][0].y;
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ arr_v : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+ @align(16)
+ arr_m : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = vec3<f32>(in[i].elements);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>) -> array<array<vec3<f32>, 9u>, 9u> {
+ var result : array<array<vec3<f32>, 9u>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+ var result : array<mat3x3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_2(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>) -> array<array<mat3x3<f32>, 9u>, 9u> {
+ var result : array<array<mat3x3<f32>, 9u>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_3(in[i]);
+ }
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_5(in : S_tint_packed_vec3) -> S {
+ var result : S;
+ result.arr_v = tint_unpack_vec3_in_composite_1(in.arr_v);
+ result.arr_m = tint_unpack_vec3_in_composite_4(in.arr_m);
+ return result;
+}
+
+fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 9u>) -> array<S, 9u> {
+ var result : array<S, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = tint_unpack_vec3_in_composite_5(in[i]);
+ }
+ return result;
+}
+
+struct S {
+ arr_v : array<array<vec3<f32>, 9>, 9>,
+ arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S_tint_packed_vec3, 9u>;
+
+fn f() {
+ var full_var : array<S, 9>;
+ var struct_var : S;
+ var outer_arr_v_var : array<array<vec3<f32>, 9>, 9>;
+ var inner_arr_v_var : array<vec3<f32>, 9>;
+ var v_var : vec3<f32>;
+ var v_element_var : f32;
+ var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9>;
+ var inner_arr_m_var : array<mat3x3<f32>, 9>;
+ var m_var : mat3x3<f32>;
+ var m_col_var : vec3<f32>;
+ var m_element_var : f32;
+ full_var = tint_unpack_vec3_in_composite_6(arr_s);
+ struct_var = tint_unpack_vec3_in_composite_5(arr_s[0]);
+ outer_arr_v_var = tint_unpack_vec3_in_composite_1(arr_s[0].arr_v);
+ inner_arr_v_var = tint_unpack_vec3_in_composite(arr_s[0].arr_v[1]);
+ v_var = vec3<f32>(arr_s[0].arr_v[1][2].elements);
+ v_element_var = arr_s[0].arr_v[1][2].elements.y;
+ outer_arr_m_var = tint_unpack_vec3_in_composite_4(arr_s[0].arr_m);
+ inner_arr_m_var = tint_unpack_vec3_in_composite_3(arr_s[0].arr_m[1]);
+ m_var = tint_unpack_vec3_in_composite_2(arr_s[0].arr_m[1][2]);
+ m_col_var = vec3<f32>(arr_s[0].arr_m[1][2][0].elements);
+ m_element_var = arr_s[0].arr_m[1][2][0].elements.y;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, RuntimeSizedArray) {
auto* src = R"(
struct S {
@@ -7817,10 +8312,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -7939,10 +8431,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
- var result : mat3x3<f32>;
- for(var i : u32; (i < 3u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
return result;
}
@@ -8048,7 +8537,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PackedVec3Test, Aliases) {
+TEST_F(PackedVec3Test, Aliases_Small) {
auto* src = R"(
alias VecArray = array<vec3<f32>, 4>;
alias MatArray = array<mat3x3<f32>, 4>;
@@ -8097,10 +8586,7 @@
}
fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
- var result : array<vec3<f32>, 4u>;
- for(var i : u32; (i < 4u); i = (i + 1)) {
- result[i] = vec3<f32>(in[i].elements);
- }
+ var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
return result;
}
@@ -8142,6 +8628,100 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, Aliases_Large) {
+ auto* src = R"(
+alias VecArray = array<vec3<f32>, 9>;
+alias MatArray = array<mat3x3<f32>, 9>;
+alias NestedArray = array<VecArray, 9>;
+
+struct S {
+ v : VecArray,
+ m : MatArray,
+ n : NestedArray,
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+@group(0) @binding(1) var<storage, read_write> arr_v : VecArray;
+@group(0) @binding(2) var<storage, read_write> arr_m : MatArray;
+@group(0) @binding(3) var<storage, read_write> arr_n : NestedArray;
+
+fn g(p : ptr<function, VecArray>) {
+}
+
+fn f() {
+ var f_arr_v : VecArray = s.v;
+ g(&f_arr_v);
+
+ arr_v = s.v;
+ arr_m[0] = s.m[0];
+ arr_n[1][2] = s.n[1][2];
+}
+)";
+
+ auto* expect =
+ R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+ @align(16)
+ elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+ @align(16)
+ v : array<tint_packed_vec3_f32_array_element, 9u>,
+ @align(16)
+ m : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
+ @align(16)
+ n : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+ var result : array<vec3<f32>, 9u>;
+ for(var i : u32; (i < 9u); i = (i + 1)) {
+ result[i] = vec3<f32>(in[i].elements);
+ }
+ return result;
+}
+
+alias VecArray = array<vec3<f32>, 9>;
+
+alias MatArray = array<mat3x3<f32>, 9>;
+
+alias NestedArray = array<VecArray, 9>;
+
+struct S {
+ v : VecArray,
+ m : MatArray,
+ n : NestedArray,
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S_tint_packed_vec3;
+
+@group(0) @binding(1) var<storage, read_write> arr_v : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(2) var<storage, read_write> arr_m : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+
+@group(0) @binding(3) var<storage, read_write> arr_n : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>;
+
+fn g(p : ptr<function, VecArray>) {
+}
+
+fn f() {
+ var f_arr_v : VecArray = tint_unpack_vec3_in_composite(s.v);
+ g(&(f_arr_v));
+ arr_v = s.v;
+ arr_m[0] = s.m[0];
+ arr_n[1][2].elements = s.n[1][2].elements;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, Vec3Bool) {
// Make sure that we don't rewrite vec3<bool> types, as the `packed_bool<n>` types are reserved
// in MSL and might not be supported everywhere.
diff --git a/src/tint/lang/spirv/reader/parser/memory_test.cc b/src/tint/lang/spirv/reader/parser/memory_test.cc
index 9f8c75f..acb3972 100644
--- a/src/tint/lang/spirv/reader/parser/memory_test.cc
+++ b/src/tint/lang/spirv/reader/parser/memory_test.cc
@@ -760,5 +760,59 @@
)");
}
+TEST_F(SpirvParserTest, StorageBufferAccessMode) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ OpDecorate %ro_var NonWritable
+ OpDecorate %ro_var DescriptorSet 1
+ OpDecorate %ro_var Binding 2
+ OpDecorate %rw_var DescriptorSet 1
+ OpDecorate %rw_var Binding 3
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %str = OpTypeStruct %u32
+ %u32_ptr = OpTypePointer StorageBuffer %u32
+ %str_ptr = OpTypePointer StorageBuffer %str
+ %ep_type = OpTypeFunction %void
+ %u32_0 = OpConstant %u32 0
+ %ro_var = OpVariable %str_ptr StorageBuffer
+ %rw_var = OpVariable %str_ptr StorageBuffer
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %ro_access = OpAccessChain %u32_ptr %ro_var %u32_0
+ %rw_access = OpAccessChain %u32_ptr %rw_var %u32_0
+ %load = OpLoad %u32 %ro_access
+ OpStore %rw_access %load
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+}
+
+%b1 = block { # root
+ %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+ %2:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 3)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+ %b2 = block {
+ %4:ptr<storage, u32, read> = access %1, 0u
+ %5:ptr<storage, u32, read_write> = access %2, 0u
+ %6:u32 = load %4
+ store %5, %6
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 6fe4eb9..cfbc8cc 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -76,6 +76,14 @@
return Failure("failed to build the internal representation of the module");
}
+ // Check for unsupported extensions.
+ for (const auto& ext : spirv_context_->extensions()) {
+ auto name = ext.GetOperand(0).AsString();
+ if (name != "SPV_KHR_storage_buffer_storage_class") {
+ return Failure("SPIR-V extension '" + name + "' is not supported");
+ }
+ }
+
{
TINT_SCOPED_ASSIGNMENT(current_block_, ir_.root_block);
EmitModuleScopeVariables();
@@ -95,10 +103,16 @@
/// @returns the Tint address space for a SPIR-V storage class
core::AddressSpace AddressSpace(spv::StorageClass sc) {
switch (sc) {
+ case spv::StorageClass::Input:
+ return core::AddressSpace::kIn;
+ case spv::StorageClass::Output:
+ return core::AddressSpace::kOut;
case spv::StorageClass::Function:
return core::AddressSpace::kFunction;
case spv::StorageClass::Private:
return core::AddressSpace::kPrivate;
+ case spv::StorageClass::StorageBuffer:
+ return core::AddressSpace::kStorage;
case spv::StorageClass::Uniform:
return core::AddressSpace::kUniform;
default:
@@ -108,10 +122,50 @@
}
}
+ /// @param b a SPIR-V BuiltIn
+ /// @returns the Tint builtin value for a SPIR-V BuiltIn decoration
+ core::BuiltinValue Builtin(spv::BuiltIn b) {
+ switch (b) {
+ case spv::BuiltIn::FragCoord:
+ return core::BuiltinValue::kPosition;
+ case spv::BuiltIn::FragDepth:
+ return core::BuiltinValue::kFragDepth;
+ case spv::BuiltIn::FrontFacing:
+ return core::BuiltinValue::kFrontFacing;
+ case spv::BuiltIn::GlobalInvocationId:
+ return core::BuiltinValue::kGlobalInvocationId;
+ case spv::BuiltIn::InstanceIndex:
+ return core::BuiltinValue::kInstanceIndex;
+ case spv::BuiltIn::LocalInvocationId:
+ return core::BuiltinValue::kLocalInvocationId;
+ case spv::BuiltIn::LocalInvocationIndex:
+ return core::BuiltinValue::kLocalInvocationIndex;
+ case spv::BuiltIn::NumWorkgroups:
+ return core::BuiltinValue::kNumWorkgroups;
+ case spv::BuiltIn::PointSize:
+ return core::BuiltinValue::kPointSize;
+ case spv::BuiltIn::Position:
+ return core::BuiltinValue::kPosition;
+ case spv::BuiltIn::SampleId:
+ return core::BuiltinValue::kSampleIndex;
+ case spv::BuiltIn::SampleMask:
+ return core::BuiltinValue::kSampleMask;
+ case spv::BuiltIn::VertexIndex:
+ return core::BuiltinValue::kVertexIndex;
+ case spv::BuiltIn::WorkgroupId:
+ return core::BuiltinValue::kWorkgroupId;
+ default:
+ TINT_UNIMPLEMENTED() << "unhandled SPIR-V BuiltIn: " << static_cast<uint32_t>(b);
+ return core::BuiltinValue::kUndefined;
+ }
+ }
+
/// @param type a SPIR-V type object
+ /// @param access_mode an optional access mode (for pointers)
/// @returns a Tint type object
- const core::type::Type* Type(const spvtools::opt::analysis::Type* type) {
- return types_.GetOrCreate(type, [&]() -> const core::type::Type* {
+ const core::type::Type* Type(const spvtools::opt::analysis::Type* type,
+ core::Access access_mode = core::Access::kUndefined) {
+ return types_.GetOrCreate(TypeKey{type, access_mode}, [&]() -> const core::type::Type* {
switch (type->kind()) {
case spvtools::opt::analysis::Type::kVoid:
return ty_.void_();
@@ -156,7 +210,7 @@
case spvtools::opt::analysis::Type::kPointer: {
auto* ptr_ty = type->AsPointer();
return ty_.ptr(AddressSpace(ptr_ty->storage_class()),
- Type(ptr_ty->pointee_type()));
+ Type(ptr_ty->pointee_type()), access_mode);
}
default:
TINT_UNIMPLEMENTED() << "unhandled SPIR-V type: " << type->str();
@@ -166,9 +220,10 @@
}
/// @param id a SPIR-V result ID for a type declaration instruction
+ /// @param access_mode an optional access mode (for pointers)
/// @returns a Tint type object
- const core::type::Type* Type(uint32_t id) {
- return Type(spirv_context_->get_type_mgr()->GetType(id));
+ const core::type::Type* Type(uint32_t id, core::Access access_mode = core::Access::kUndefined) {
+ return Type(spirv_context_->get_type_mgr()->GetType(id), access_mode);
}
/// @param arr_ty a SPIR-V array object
@@ -209,6 +264,15 @@
uint32_t align = std::max<uint32_t>(member_ty->Align(), 1u);
uint32_t offset = tint::RoundUp(align, current_size);
core::type::StructMemberAttributes attributes;
+ auto interpolation = [&]() -> core::Interpolation& {
+ // Create the interpolation field with the default values on first call.
+ if (!attributes.interpolation.has_value()) {
+ attributes.interpolation =
+ core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCenter};
+ }
+ return attributes.interpolation.value();
+ };
// Handle member decorations that affect layout or attributes.
if (struct_ty->element_decorations().count(i)) {
@@ -217,6 +281,28 @@
case spv::Decoration::Offset:
offset = deco[1];
break;
+ case spv::Decoration::BuiltIn:
+ attributes.builtin = Builtin(spv::BuiltIn(deco[1]));
+ break;
+ case spv::Decoration::Invariant:
+ attributes.invariant = true;
+ break;
+ case spv::Decoration::Location:
+ attributes.location = deco[1];
+ break;
+ case spv::Decoration::NoPerspective:
+ interpolation().type = core::InterpolationType::kLinear;
+ break;
+ case spv::Decoration::Flat:
+ interpolation().type = core::InterpolationType::kFlat;
+ break;
+ case spv::Decoration::Centroid:
+ interpolation().sampling = core::InterpolationSampling::kCentroid;
+ break;
+ case spv::Decoration::Sample:
+ interpolation().sampling = core::InterpolationSampling::kSample;
+ break;
+
default:
TINT_UNIMPLEMENTED() << "unhandled member decoration: " << deco[0];
break;
@@ -459,7 +545,14 @@
indices.Push(Value(inst.GetSingleWordOperand(i)));
}
auto* base = Value(inst.GetSingleWordOperand(2));
- auto* access = b_.Access(Type(inst.type_id()), base, std::move(indices));
+
+ // Propagate the access mode of the base object.
+ auto access_mode = core::Access::kUndefined;
+ if (auto* ptr = base->Type()->As<core::type::Pointer>()) {
+ access_mode = ptr->Access();
+ }
+
+ auto* access = b_.Access(Type(inst.type_id(), access_mode), base, std::move(indices));
Emit(access, inst.result_id());
}
@@ -496,19 +589,25 @@
/// @param inst the SPIR-V instruction for OpVariable
void EmitVar(const spvtools::opt::Instruction& inst) {
- auto* var = b_.Var(Type(inst.type_id())->As<core::type::Pointer>());
- if (inst.NumOperands() > 3) {
- var->SetInitializer(Value(inst.GetSingleWordOperand(3)));
- }
-
// Handle decorations.
std::optional<uint32_t> group;
std::optional<uint32_t> binding;
+ core::Access access_mode = core::Access::kUndefined;
+ core::ir::IOAttributes io_attributes;
+ auto interpolation = [&]() -> core::Interpolation& {
+ // Create the interpolation field with the default values on first call.
+ if (!io_attributes.interpolation.has_value()) {
+ io_attributes.interpolation = core::Interpolation{
+ core::InterpolationType::kPerspective, core::InterpolationSampling::kCenter};
+ }
+ return io_attributes.interpolation.value();
+ };
for (auto* deco :
spirv_context_->get_decoration_mgr()->GetDecorationsFor(inst.result_id(), false)) {
auto d = deco->GetSingleWordOperand(1);
switch (spv::Decoration(d)) {
case spv::Decoration::NonWritable:
+ access_mode = core::Access::kRead;
break;
case spv::Decoration::DescriptorSet:
group = deco->GetSingleWordOperand(2);
@@ -516,20 +615,70 @@
case spv::Decoration::Binding:
binding = deco->GetSingleWordOperand(2);
break;
+ case spv::Decoration::BuiltIn:
+ io_attributes.builtin = Builtin(spv::BuiltIn(deco->GetSingleWordOperand(2)));
+ break;
+ case spv::Decoration::Invariant:
+ io_attributes.invariant = true;
+ break;
+ case spv::Decoration::Location:
+ io_attributes.location = deco->GetSingleWordOperand(2);
+ break;
+ case spv::Decoration::NoPerspective:
+ interpolation().type = core::InterpolationType::kLinear;
+ break;
+ case spv::Decoration::Flat:
+ interpolation().type = core::InterpolationType::kFlat;
+ break;
+ case spv::Decoration::Centroid:
+ interpolation().sampling = core::InterpolationSampling::kCentroid;
+ break;
+ case spv::Decoration::Sample:
+ interpolation().sampling = core::InterpolationSampling::kSample;
+ break;
default:
TINT_UNIMPLEMENTED() << "unhandled decoration " << d;
break;
}
}
+
+ auto* var = b_.Var(Type(inst.type_id(), access_mode)->As<core::type::Pointer>());
+ if (inst.NumOperands() > 3) {
+ var->SetInitializer(Value(inst.GetSingleWordOperand(3)));
+ }
+
if (group || binding) {
TINT_ASSERT(group && binding);
var->SetBindingPoint(group.value(), binding.value());
}
+ var->SetAttributes(std::move(io_attributes));
Emit(var, inst.result_id());
}
private:
+ /// TypeKey describes a SPIR-V type with an access mode.
+ struct TypeKey {
+ /// The SPIR-V type object.
+ const spvtools::opt::analysis::Type* type;
+ /// The access mode.
+ core::Access access_mode;
+
+ // Equality operator for TypeKey.
+ bool operator==(const TypeKey& other) const {
+ return type == other.type && access_mode == other.access_mode;
+ }
+
+ /// Hasher provides a hash function for the TypeKey.
+ struct Hasher {
+ /// @param tk the TypeKey to create a hash for
+ /// @return the hash value
+ inline std::size_t operator()(const TypeKey& tk) const {
+ return HashCombine(Hash(tk.type), tk.access_mode);
+ }
+ };
+ };
+
/// The generated IR module.
core::ir::Module ir_;
/// The Tint IR builder.
@@ -541,8 +690,8 @@
core::ir::Function* current_function_ = nullptr;
/// The Tint IR block that is currently being emitted.
core::ir::Block* current_block_ = nullptr;
- /// A map from a SPIR-V type declaration result ID to the corresponding Tint type object.
- Hashmap<const spvtools::opt::analysis::Type*, const core::type::Type*, 16> types_;
+ /// A map from a SPIR-V type declaration to the corresponding Tint type object.
+ Hashmap<TypeKey, const core::type::Type*, 16, TypeKey::Hasher> types_;
/// A map from a SPIR-V function definition result ID to the corresponding Tint function object.
Hashmap<uint32_t, core::ir::Function*, 8> functions_;
/// A map from a SPIR-V result ID to the corresponding Tint value object.
diff --git a/src/tint/lang/spirv/reader/parser/struct_test.cc b/src/tint/lang/spirv/reader/parser/struct_test.cc
index afbaa39..bedb4e2 100644
--- a/src/tint/lang/spirv/reader/parser/struct_test.cc
+++ b/src/tint/lang/spirv/reader/parser/struct_test.cc
@@ -229,4 +229,132 @@
)");
}
+TEST_F(SpirvParserTest, Struct_Builtin) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ OpMemberDecorate %str 0 BuiltIn Position
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %str = OpTypeStruct %vec4f
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %str
+ %var = OpVariable %_ptr_Output Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @builtin(position)
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Struct_Builtin_WithInvariant) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ OpMemberDecorate %str 0 BuiltIn Position
+ OpMemberDecorate %str 0 Invariant
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %str = OpTypeStruct %vec4f
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %str
+ %var = OpVariable %_ptr_Output Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @invariant, @builtin(position)
+}
+)");
+}
+
+struct LocationCase {
+ std::string spirv_decorations;
+ std::string ir;
+};
+
+using LocationStructTest = SpirvParserTestWithParam<LocationCase>;
+
+TEST_P(LocationStructTest, MemberDecorations) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main"
+ OpExecutionMode %main OriginUpperLeft
+ )" + params.spirv_decorations +
+ R"(
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %str = OpTypeStruct %vec4f
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Input = OpTypePointer Input %str
+ %var = OpVariable %_ptr_Input Input
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ SpirvParser,
+ LocationStructTest,
+ testing::Values(
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 1 ",
+ "tint_symbol:vec4<f32> @offset(0), @location(1)",
+ },
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 2 "
+ "OpMemberDecorate %str 0 NoPerspective ",
+ "tint_symbol:vec4<f32> @offset(0), @location(2), @interpolate(linear, center)",
+ },
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 3 "
+ "OpMemberDecorate %str 0 Flat ",
+ "tint_symbol:vec4<f32> @offset(0), @location(3), @interpolate(flat, center)",
+ },
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 4 "
+ "OpMemberDecorate %str 0 Centroid ",
+ "tint_symbol:vec4<f32> @offset(0), @location(4), @interpolate(perspective, centroid)",
+ },
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 5 "
+ "OpMemberDecorate %str 0 Sample ",
+ "tint_symbol:vec4<f32> @offset(0), @location(5), @interpolate(perspective, sample)",
+ },
+ LocationCase{
+ "OpMemberDecorate %str 0 Location 6 "
+ "OpMemberDecorate %str 0 NoPerspective "
+ "OpMemberDecorate %str 0 Centroid ",
+ "tint_symbol:vec4<f32> @offset(0), @location(6), @interpolate(linear, centroid)",
+ }));
+
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 88a6347..f524b23 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -142,6 +142,129 @@
)");
}
+TEST_F(SpirvParserTest, StorageVar_ReadOnly) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %1 "main"
+ OpExecutionMode %1 LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ OpDecorate %6 NonWritable
+ OpDecorate %6 DescriptorSet 1
+ OpDecorate %6 Binding 2
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+ %5 = OpTypeFunction %void
+ %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+ %1 = OpFunction %void None %5
+ %7 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+}
+
+%b1 = block { # root
+ %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, StorageVar_ReadWrite) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %1 "main"
+ OpExecutionMode %1 LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ OpDecorate %6 DescriptorSet 1
+ OpDecorate %6 Binding 2
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+ %5 = OpTypeFunction %void
+ %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+ %1 = OpFunction %void None %5
+ %7 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+}
+
+%b1 = block { # root
+ %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 2)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, StorageVar_ReadOnly_And_ReadWrite) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %1 "main"
+ OpExecutionMode %1 LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ OpDecorate %6 NonWritable
+ OpDecorate %6 DescriptorSet 1
+ OpDecorate %6 Binding 2
+ OpDecorate %7 DescriptorSet 1
+ OpDecorate %7 Binding 3
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+ %5 = OpTypeFunction %void
+ %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+ %7 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+ %1 = OpFunction %void None %5
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+}
+
+%b1 = block { # root
+ %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+ %2:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 3)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
TEST_F(SpirvParserTest, UniformVar) {
EXPECT_IR(R"(
OpCapability Shader
@@ -181,5 +304,278 @@
)");
}
+struct BuiltinCase {
+ std::string spirv_type;
+ std::string spirv_builtin;
+ std::string ir;
+};
+std::string PrintBuiltinCase(testing::TestParamInfo<BuiltinCase> bc) {
+ return bc.param.spirv_builtin;
+}
+
+using BuiltinInputTest = SpirvParserTestWithParam<BuiltinCase>;
+
+TEST_P(BuiltinInputTest, Enum) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %var BuiltIn )" +
+ params.spirv_builtin + R"(
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %vec3u = OpTypeVector %u32 3
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %u32_1 = OpConstant %u32 1
+ %arr_u32_1 = OpTypeArray %u32 %u32_1
+ %fn_type = OpTypeFunction %void
+
+ %_ptr_Input = OpTypePointer Input %)" +
+ params.spirv_type + R"(
+ %var = OpVariable %_ptr_Input Input
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ SpirvParser,
+ BuiltinInputTest,
+ testing::Values(
+ BuiltinCase{
+ "vec4f",
+ "FragCoord",
+ "%1:ptr<__in, vec4<f32>, read> = var @builtin(position)",
+ },
+ BuiltinCase{
+ "bool",
+ "FrontFacing",
+ "%1:ptr<__in, bool, read> = var @builtin(front_facing)",
+ },
+ BuiltinCase{
+ "vec3u",
+ "GlobalInvocationId",
+ "%1:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)",
+ },
+ BuiltinCase{
+ "u32",
+ "InstanceIndex",
+ "%1:ptr<__in, u32, read> = var @builtin(instance_index)",
+ },
+ BuiltinCase{
+ "vec3u",
+ "LocalInvocationId",
+ "%1:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)",
+ },
+ BuiltinCase{
+ "u32",
+ "LocalInvocationIndex",
+ "%1:ptr<__in, u32, read> = var @builtin(local_invocation_index)",
+ },
+ BuiltinCase{
+ "vec3u",
+ "NumWorkgroups",
+ "%1:ptr<__in, vec3<u32>, read> = var @builtin(num_workgroups)",
+ },
+ BuiltinCase{
+ "u32",
+ "SampleId",
+ "%1:ptr<__in, u32, read> = var @builtin(sample_index)",
+ },
+ BuiltinCase{
+ "arr_u32_1",
+ "SampleMask",
+ "%1:ptr<__in, array<u32, 1>, read> = var @builtin(sample_mask)",
+ },
+ BuiltinCase{
+ "u32",
+ "VertexIndex",
+ "%1:ptr<__in, u32, read> = var @builtin(vertex_index)",
+ },
+ BuiltinCase{
+ "vec3u",
+ "WorkgroupId",
+ "%1:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)",
+ }),
+ PrintBuiltinCase);
+
+using BuiltinOutputTest = SpirvParserTestWithParam<BuiltinCase>;
+
+TEST_P(BuiltinOutputTest, Enum) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %var BuiltIn )" +
+ params.spirv_builtin + R"(
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %u32_1 = OpConstant %u32 1
+ %arr_u32_1 = OpTypeArray %u32 %u32_1
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %)" +
+ params.spirv_type + R"(
+ %var = OpVariable %_ptr_Output Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ SpirvParser,
+ BuiltinOutputTest,
+ testing::Values(
+ BuiltinCase{
+ "f32",
+ "PointSize",
+ "%1:ptr<__out, f32, read_write> = var @builtin(__point_size)",
+ },
+ BuiltinCase{
+ "vec4f",
+ "Position",
+ "%1:ptr<__out, vec4<f32>, read_write> = var @builtin(position)",
+ },
+ BuiltinCase{
+ "arr_u32_1",
+ "SampleMask",
+ "%1:ptr<__out, array<u32, 1>, read_write> = var @builtin(sample_mask)",
+ }),
+ PrintBuiltinCase);
+
+TEST_F(SpirvParserTest, Invariant_OnVariable) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %var BuiltIn Position
+ OpDecorate %var Invariant
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %vec4f
+ %var = OpVariable %_ptr_Output Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ "%1:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)");
+}
+
+struct LocationCase {
+ std::string spirv_decorations;
+ std::string ir;
+};
+
+using LocationVarTest = SpirvParserTestWithParam<LocationCase>;
+
+TEST_P(LocationVarTest, Input) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ )" + params.spirv_decorations +
+ R"(
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Input = OpTypePointer Input %vec4f
+ %var = OpVariable %_ptr_Input Input
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ "%1:ptr<__in, vec4<f32>, read> = " + params.ir);
+}
+
+TEST_P(LocationVarTest, Output) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ )" + params.spirv_decorations +
+ R"(
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %vec4f
+ %var = OpVariable %_ptr_Output Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ "%1:ptr<__out, vec4<f32>, read_write> = " + params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+ LocationVarTest,
+ testing::Values(
+ LocationCase{
+ "OpDecorate %var Location 1 ",
+ "var @location(1)",
+ },
+ LocationCase{
+ "OpDecorate %var Location 2 "
+ "OpDecorate %var NoPerspective ",
+ "var @location(2) @interpolate(linear, center)",
+ },
+ LocationCase{
+ "OpDecorate %var Location 3 "
+ "OpDecorate %var Flat ",
+ "var @location(3) @interpolate(flat, center)",
+ },
+ LocationCase{
+ "OpDecorate %var Location 4 "
+ "OpDecorate %var Centroid ",
+ "var @location(4) @interpolate(perspective, centroid)",
+ },
+ LocationCase{
+ "OpDecorate %var Location 5 "
+ "OpDecorate %var Sample ",
+ "var @location(5) @interpolate(perspective, sample)",
+ },
+ LocationCase{
+ "OpDecorate %var Location 6 "
+ "OpDecorate %var NoPerspective "
+ "OpDecorate %var Centroid ",
+ "var @location(6) @interpolate(linear, centroid)",
+ }));
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index 1a086a8..5313aaf 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -68,6 +68,25 @@
}
};
+TEST_F(SpirvReaderTest, UnsupportedExtension) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_variable_pointers"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_NE(got, Success);
+ EXPECT_EQ(got.Failure().reason.str(),
+ "error: SPIR-V extension 'SPV_KHR_variable_pointers' is not supported");
+}
+
TEST_F(SpirvReaderTest, Load_VectorComponent) {
auto got = Run(R"(
OpCapability Shader
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index 630ef0f..d7fe564 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -27,7 +27,7 @@
#include "src/tint/lang/wgsl/inspector/inspector.h"
-#include <limits>
+#include <unordered_set>
#include <utility>
#include "src/tint/lang/core/builtin_value.h"
@@ -1009,4 +1009,139 @@
}
}
+std::vector<Inspector::LevelSampleInfo> Inspector::GetTextureQueries(const std::string& ep_name) {
+ std::vector<LevelSampleInfo> res;
+
+ std::unordered_set<BindingPoint> seen = {};
+
+ auto sample_type_for_call_and_type = [](wgsl::BuiltinFn builtin, const core::type::Type* ty) {
+ if (builtin == wgsl::BuiltinFn::kTextureNumLevels) {
+ return TextureQueryType::kTextureNumLevels;
+ }
+ if (builtin == wgsl::BuiltinFn::kTextureLoad) {
+ if (!ty->UnwrapRef()
+ ->IsAnyOf<core::type::MultisampledTexture,
+ core::type::DepthMultisampledTexture>()) {
+ return TextureQueryType::kTextureNumLevels;
+ }
+ }
+
+ return TextureQueryType::kTextureNumSamples;
+ };
+
+ Hashmap<const sem::Function*, Hashmap<const ast::Parameter*, TextureQueryType, 4>, 8>
+ fn_to_data;
+
+ auto record_function_param = [&fn_to_data](const sem::Function* func,
+ const ast::Parameter* param, TextureQueryType type) {
+ auto& param_to_type = *fn_to_data.GetOrZero(func);
+
+ auto entry = param_to_type.Get(param);
+ if (entry.has_value()) {
+ return;
+ }
+
+ param_to_type.Add(param, type);
+ };
+
+ auto save_if_needed = [&res, &seen](const sem::GlobalVariable* global, TextureQueryType type) {
+ auto binding = global->Attributes().binding_point.value();
+ if (seen.insert(binding).second) {
+ res.emplace_back(LevelSampleInfo{type, binding.group, binding.binding});
+ }
+ };
+
+ auto& sem = program_.Sem();
+
+ const auto* ep = FindEntryPointByName(ep_name);
+ if (!ep) {
+ return {};
+ }
+
+ // This works in dependency order such that we'll see the texture call first and can record
+ // any function parameter information and then as we walk up the function chain we can look
+ // the call data.
+ for (auto* fn_decl : sem.Module()->DependencyOrderedDeclarations()) {
+ auto* fn = sem.Get<sem::Function>(fn_decl);
+ if (!fn) {
+ continue;
+ }
+
+ // This is an entrypoint, make sure it's the requested entry point
+ if (fn->Declaration()->IsEntryPoint()) {
+ if (fn->Declaration() != ep) {
+ continue;
+ }
+ } else {
+ // Not an entry point, make sure it was called from the requested entry point
+ if (!fn->HasAncestorEntryPoint(ep->name->symbol)) {
+ continue;
+ }
+ }
+
+ for (auto* call : fn->DirectCalls()) {
+ // Builtin function call, record the texture information. If the used texture maps
+ // back up to a function parameter just store the type of the call and we'll track the
+ // function callback up in the `sem::Function` branch.
+ tint::Switch(
+ call->Target(),
+ [&](const sem::BuiltinFn* builtin) {
+ if (builtin->Fn() != wgsl::BuiltinFn::kTextureNumLevels &&
+ builtin->Fn() != wgsl::BuiltinFn::kTextureNumSamples &&
+ builtin->Fn() != wgsl::BuiltinFn::kTextureLoad) {
+ return;
+ }
+
+ auto* texture_expr = call->Declaration()->args[0];
+ auto* texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
+ TINT_ASSERT(texture_sem);
+
+ auto type = sample_type_for_call_and_type(builtin->Fn(), texture_sem->Type());
+
+ tint::Switch(
+ texture_sem, //
+ [&](const sem::GlobalVariable* global) { save_if_needed(global, type); },
+ [&](const sem::Parameter* param) {
+ record_function_param(fn, param->Declaration(), type);
+ },
+ TINT_ICE_ON_NO_MATCH);
+ },
+ [&](const sem::Function* func) {
+ // A function call, check to see if any params needed to be tracked back to a
+ // global texture.
+
+ auto param_to_type = fn_to_data.Find(func);
+ if (!param_to_type) {
+ return;
+ }
+ TINT_ASSERT(call->Arguments().Length() == func->Declaration()->params.Length());
+
+ for (size_t i = 0; i < call->Arguments().Length(); i++) {
+ auto param = func->Declaration()->params[i];
+
+ // Determine if this had a texture we cared about
+ auto type = param_to_type->Get(param);
+ if (!type.has_value()) {
+ continue;
+ }
+
+ auto* arg = call->Arguments()[i];
+ auto* texture_sem = arg->RootIdentifier();
+
+ tint::Switch(
+ texture_sem,
+ [&](const sem::GlobalVariable* global) {
+ save_if_needed(global, type.value());
+ },
+ [&](const sem::Parameter* p) {
+ record_function_param(fn, p->Declaration(), type.value());
+ },
+ TINT_ICE_ON_NO_MATCH);
+ }
+ });
+ }
+ }
+ return res;
+}
+
} // namespace tint::inspector
diff --git a/src/tint/lang/wgsl/inspector/inspector.h b/src/tint/lang/wgsl/inspector/inspector.h
index 14c5ef6..6604832 100644
--- a/src/tint/lang/wgsl/inspector/inspector.h
+++ b/src/tint/lang/wgsl/inspector/inspector.h
@@ -156,6 +156,31 @@
/// extension.
std::vector<std::pair<std::string, Source>> GetEnableDirectives();
+ /// The information needed to be supplied.
+ enum class TextureQueryType : uint8_t {
+ /// Texture Num Levels
+ kTextureNumLevels,
+ /// Texture Num Samples
+ kTextureNumSamples,
+ };
+ /// Information on level and sample calls by a given texture binding point
+ struct LevelSampleInfo {
+ /// The type of function
+ TextureQueryType type = TextureQueryType::kTextureNumLevels;
+ /// The group number
+ uint32_t group = 0;
+ /// The binding number
+ uint32_t binding = 0;
+ };
+
+ /// @param ep the entry point ot get the information for
+ /// @returns a vector of information for textures which call textureNumLevels and
+ /// textureNumSamples for backends which require additional support for those methods. Each
+ /// binding point will only be returned once regardless of the number of calls made. The
+ /// texture types for `textureNumSamples` is disjoint from the texture types in
+ /// `textureNumLevels` so the binding point will always be one or the other.
+ std::vector<LevelSampleInfo> GetTextureQueries(const std::string& ep);
+
private:
const Program& program_;
diag::List diagnostics_;
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index c33f950..76d514e 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -3701,5 +3701,286 @@
inspector.GetSamplerTextureUses("main");
}
+class InspectorTextureTest : public InspectorRunner, public testing::Test {};
+
+TEST_F(InspectorTextureTest, TextureLevelInEP) {
+ std::string shader = R"(
+@group(2) @binding(3) var myTexture: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num = textureNumLevels(myTexture);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLevelInEPNoDups) {
+ std::string shader = R"(
+@group(0) @binding(0) var myTexture: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureNumLevels(myTexture);
+ let num2 = textureNumLevels(myTexture);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+}
+
+TEST_F(InspectorTextureTest, TextureLevelInEPMultiple) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(2) var tex2: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureNumLevels(tex1);
+ let num2 = textureNumLevels(tex2);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(2u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[1].type);
+ EXPECT_EQ(1u, info[1].group);
+ EXPECT_EQ(2u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEP) {
+ std::string shader = R"(
+@group(2) @binding(3) var myTexture: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num = textureNumSamples(myTexture);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEPNoDups) {
+ std::string shader = R"(
+@group(0) @binding(0) var myTexture: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureNumSamples(myTexture);
+ let num2 = textureNumSamples(myTexture);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEPMultiple) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_multisampled_2d<f32>;
+@group(1) @binding(2) var tex2: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureNumSamples(tex1);
+ let num2 = textureNumSamples(tex2);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(2u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+ EXPECT_EQ(1u, info[1].group);
+ EXPECT_EQ(2u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadInEP) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureLoad(tex1, vec2(0, 0), 0);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadMultisampledInEP) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureLoad(tex1, vec2(0, 0), 0);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(1u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadMultipleInEP) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(4) var tex2: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let num1 = textureLoad(tex1, vec2(0, 0), 0);
+ let num2 = textureLoad(tex2, vec2(0, 0), 0);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(2u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+ EXPECT_EQ(1u, info[1].group);
+ EXPECT_EQ(4u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureInSubfunction) {
+ std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(4) var tex2: texture_multisampled_2d<f32>;
+@group(1) @binding(3) var tex3: texture_2d<f32>;
+
+fn b(tx1: texture_2d<f32>, tx2: texture_multisampled_2d<f32>, tx3: texture_2d<f32>, tx4: texture_2d<f32>) {
+ let v1 = textureNumLevels(tx1);
+ let v2 = textureNumSamples(tx2);
+ let v3 = textureLoad(tx3, vec2(0, 0), 0);
+ let v4 = textureNumLevels(tx4);
+}
+
+fn a(tx1: texture_2d<f32>, tx2: texture_multisampled_2d<f32>, tx3: texture_2d<f32>) {
+ b(tx1, tx2, tx3, tx1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ a(tex1, tex2, tex3);
+})";
+
+ Inspector& inspector = Initialize(shader);
+ auto info = inspector.GetTextureQueries("main");
+
+ ASSERT_EQ(3u, info.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+ EXPECT_EQ(2u, info[0].group);
+ EXPECT_EQ(3u, info[0].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+ EXPECT_EQ(1u, info[1].group);
+ EXPECT_EQ(4u, info[1].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[2].type);
+ EXPECT_EQ(1u, info[2].group);
+ EXPECT_EQ(3u, info[2].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureMultipleEPs) {
+ std::string shader = R"(
+@group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>;
+@group(0) @binding(1) var tex1 : texture_2d_array<f32>;
+@group(0) @binding(4) var tex2 : texture_multisampled_2d<f32>;
+@group(1) @binding(3) var tex3 : texture_2d_array<f32>;
+
+@compute @workgroup_size(1, 1, 1) fn main1() {
+ dstBuf[0] = textureNumLayers(tex1);
+ dstBuf[1] = textureNumLevels(tex1);
+ dstBuf[2] = textureNumSamples(tex2);
+ dstBuf[3] = textureNumLevels(tex3);
+}
+
+@compute @workgroup_size(1, 1, 1) fn main2() {
+ dstBuf[0] = textureNumLayers(tex1);
+ dstBuf[1] = textureNumLevels(tex1);
+ dstBuf[2] = textureNumSamples(tex2);
+}
+ )";
+ Inspector& inspector = Initialize(shader);
+ {
+ auto info1 = inspector.GetTextureQueries("main1");
+ ASSERT_EQ(3u, info1.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info1[0].type);
+ EXPECT_EQ(0u, info1[0].group);
+ EXPECT_EQ(1u, info1[0].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info1[1].type);
+ EXPECT_EQ(0u, info1[1].group);
+ EXPECT_EQ(4u, info1[1].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info1[2].type);
+ EXPECT_EQ(1u, info1[2].group);
+ EXPECT_EQ(3u, info1[2].binding);
+ }
+ {
+ auto info2 = inspector.GetTextureQueries("main2");
+ ASSERT_EQ(2u, info2.size());
+
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info2[0].type);
+ EXPECT_EQ(0u, info2[0].group);
+ EXPECT_EQ(1u, info2[0].binding);
+ EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info2[1].type);
+ EXPECT_EQ(0u, info2[1].group);
+ EXPECT_EQ(4u, info2[1].binding);
+ }
+}
+
} // namespace
+
+static std::ostream& operator<<(std::ostream& out, const Inspector::TextureQueryType& ty) {
+ switch (ty) {
+ case Inspector::TextureQueryType::kTextureNumLevels:
+ out << "textureNumLevels";
+ break;
+ case Inspector::TextureQueryType::kTextureNumSamples:
+ out << "textureNumSamples";
+ break;
+ }
+ return out;
+}
+
} // namespace tint::inspector
