[spirv] Fix dynamic indexes into constant arrays
When a constant array was dynamically indexed from multiple functions,
the VarForDynamicIndex transform was adding a function-scope copy of
the array to one function and then trying to use it from the other
functions. This CL changes the transform to use module-scope variables
for constant array indexing instead.
A similar problem exists for non-constant arrays that are indexed from
multiple blocks. This is fixed here by moving the function-scope
declaration to the earliest point that is immediately after the
definition of the array value. A new helper is added to the IR builder
to determine the correct insertion point depending on the kind of
value we are inserting after (instruction result, function parameter,
or block parameter).
Fixed: tint:2237
Change-Id: Ib7d9c3ae7f83939c0270c8da6ce6c29994b88de5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/185525
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index c705d0c..99e94d1 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -219,6 +219,37 @@
cb();
}
+ /// Calls @p cb with the builder inserting after @p val
+ /// @param val the insertion point for new instructions
+ /// @param cb the function to call with the builder inserting new instructions after @p val
+ template <typename FUNCTION>
+ void InsertAfter(ir::Value* val, FUNCTION&& cb) {
+ tint::Switch(
+ val,
+ [&](core::ir::InstructionResult* result) {
+ const TINT_SCOPED_ASSIGNMENT(insertion_point_,
+ InsertionPoints::InsertAfter{result->Instruction()});
+ cb();
+ },
+ [&](core::ir::FunctionParam* param) {
+ auto* body = param->Function()->Block();
+ if (body->IsEmpty()) {
+ Append(body, cb);
+ } else {
+ InsertBefore(body->Front(), cb);
+ }
+ },
+ [&](core::ir::BlockParam* param) {
+ auto* block = param->Block();
+ if (block->IsEmpty()) {
+ Append(block, cb);
+ } else {
+ InsertBefore(block->Front(), cb);
+ }
+ },
+ TINT_ICE_ON_NO_MATCH);
+ }
+
/// Adds and returns the instruction @p instruction to the current insertion point. If there
/// is no current insertion point set, then @p instruction is just returned.
/// @param instruction the instruction to append
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
index 901b4bb..98b3067 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
@@ -134,7 +134,7 @@
}
// Replace each access instruction that we recorded.
- Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_local;
+ Hashmap<core::ir::Value*, core::ir::Value*, 4> object_to_var;
Hashmap<PartialAccess, core::ir::Value*, 4> source_object_to_value;
for (const auto& to_replace : worklist) {
auto* access = to_replace.access;
@@ -145,24 +145,58 @@
if (to_replace.first_dynamic_index > 0) {
PartialAccess partial_access = {
access->Object(), access->Indices().Truncate(to_replace.first_dynamic_index)};
- source_object = source_object_to_value.GetOrAdd(partial_access, [&] {
- auto* intermediate_source = builder.Access(to_replace.dynamic_index_source_type,
- source_object, partial_access.indices);
- intermediate_source->InsertBefore(access);
- return intermediate_source->Result(0);
- });
+ source_object =
+ source_object_to_value.GetOrAdd(partial_access, [&]() -> core::ir::Value* {
+ // If the source is a constant, then the partial access will also produce a
+ // constant. Extract the constant::Value and use that as the new source object.
+ if (source_object->Is<core::ir::Constant>()) {
+ for (const auto& i : partial_access.indices) {
+ auto idx = i->As<core::ir::Constant>()->Value()->ValueAs<uint32_t>();
+ source_object = builder.Constant(
+ source_object->As<core::ir::Constant>()->Value()->Index(idx));
+ }
+ return source_object;
+ }
+
+ // Extract a non-constant intermediate source using an access instruction that
+ // we insert immediately after the definition of the root source object.
+ auto* intermediate_source =
+ builder.Access(to_replace.dynamic_index_source_type, source_object,
+ partial_access.indices);
+ builder.InsertAfter(source_object,
+ [&] { builder.Append(intermediate_source); });
+ return intermediate_source->Result(0);
+ });
}
- // Declare a local variable and copy the source object to it.
- auto* local = object_to_local.GetOrAdd(source_object, [&] {
- auto* decl = builder.Var(ir.Types().ptr(
- core::AddressSpace::kFunction, source_object->Type(), core::Access::kReadWrite));
+ // Declare a variable and copy the source object to it.
+ auto* var = object_to_var.GetOrAdd(source_object, [&] {
+ // If the source object is a constant we use a module-scope variable, as it could be
+ // indexed by multiple functions. Otherwise, we declare a function-scope variable
+ // immediately after the definition of the source object.
+ core::ir::Var* decl = nullptr;
+ if (source_object->Is<core::ir::Constant>()) {
+ decl = builder.Var(ir.Types().ptr(core::AddressSpace::kPrivate,
+ source_object->Type(), core::Access::kReadWrite));
+ ir.root_block->Append(decl);
+ } else {
+ builder.InsertAfter(source_object, [&] {
+ decl = builder.Var(ir.Types().ptr(core::AddressSpace::kFunction,
+ source_object->Type(),
+ core::Access::kReadWrite));
+
+ // If we ever support value declarations at module-scope, we will need to modify
+ // the partial access logic above since `access` instructions cannot be used in
+ // the root block.
+ TINT_ASSERT(decl->Block() != ir.root_block);
+ });
+ }
+
decl->SetInitializer(source_object);
- decl->InsertBefore(access);
return decl->Result(0);
});
- // Create a new access instruction using the local variable as the source.
+ // Create a new access instruction using the new variable as the source.
Vector<core::ir::Value*, 4> indices{
access->Indices().Offset(to_replace.first_dynamic_index)};
const core::type::Type* access_type = access->Result(0)->Type();
@@ -178,9 +212,9 @@
vector_index = indices.Pop();
}
+ auto addrspace = var->Type()->As<core::type::Pointer>()->AddressSpace();
core::ir::Instruction* new_access = builder.Access(
- ir.Types().ptr(core::AddressSpace::kFunction, access_type, core::Access::kReadWrite),
- local, indices);
+ ir.Types().ptr(addrspace, access_type, core::Access::kReadWrite), var, indices);
new_access->InsertBefore(access);
core::ir::Instruction* load = nullptr;
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
index e4adae9..657488b 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index_test.cc
@@ -430,5 +430,449 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToFuncParam_FromDifferentBlocks) {
+ auto* arr = b.FunctionParam(ty.array<i32, 4>());
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({arr, cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_b));
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %6:i32 = access %2, %4
+ ret %6
+ }
+ %b3 = block { # false
+ %7:i32 = access %2, %5
+ ret %7
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:array<i32, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ %6:ptr<function, array<i32, 4>, read_write> = var, %2
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %7:ptr<function, i32, read_write> = access %6, %4
+ %8:i32 = load %7
+ ret %8
+ }
+ %b3 = block { # false
+ %9:ptr<function, i32, read_write> = access %6, %5
+ %10:i32 = load %9
+ ret %10
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToFuncParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+ auto* arr = b.FunctionParam(ty.array(ty.array<i32, 4>(), 4));
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({arr, cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %6:i32 = access %2, 0u, %4
+ ret %6
+ }
+ %b3 = block { # false
+ %7:i32 = access %2, 0u, %5
+ ret %7
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:array<array<i32, 4>, 4>, %3:bool, %4:i32, %5:i32):i32 -> %b1 {
+ %b1 = block {
+ %6:array<i32, 4> = access %2, 0u
+ %7:ptr<function, array<i32, 4>, read_write> = var, %6
+ if %3 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ %8:ptr<function, i32, read_write> = access %7, %4
+ %9:i32 = load %8
+ ret %9
+ }
+ %b3 = block { # false
+ %10:ptr<function, i32, read_write> = access %7, %5
+ %11:i32 = load %10
+ ret %11
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToBlockParam_FromDifferentBlocks) {
+ auto* arr = b.BlockParam(ty.array<i32, 4>());
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ loop->Body()->SetParams({arr});
+ b.Append(loop->Body(), [&] {
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, idx_b));
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<i32, 4>) { # body
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %6:i32 = access %5:array<i32, 4>, %3
+ ret %6
+ }
+ %b4 = block { # false
+ %7:i32 = access %5:array<i32, 4>, %4
+ ret %7
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<i32, 4>) { # body
+ %6:ptr<function, array<i32, 4>, read_write> = var, %5:array<i32, 4>
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %7:ptr<function, i32, read_write> = access %6, %3
+ %8:i32 = load %7
+ ret %8
+ }
+ %b4 = block { # false
+ %9:ptr<function, i32, read_write> = access %6, %4
+ %10:i32 = load %9
+ ret %10
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToBlockParam_FromDifferentBlocks_WithLeadingConstantIndex) {
+ auto* arr = b.BlockParam(ty.array(ty.array<i32, 4>(), 4));
+ auto* cond = b.FunctionParam(ty.bool_());
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func = b.Function("func", ty.i32());
+ func->SetParams({cond, idx_a, idx_b});
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ loop->Body()->SetParams({arr});
+ b.Append(loop->Body(), [&] {
+ auto* if_ = b.If(cond);
+ b.Append(if_->True(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+ b.Append(if_->False(), [&] { //
+ b.Return(func, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<array<i32, 4>, 4>) { # body
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %6:i32 = access %5:array<array<i32, 4>, 4>, 0u, %3
+ ret %6
+ }
+ %b4 = block { # false
+ %7:i32 = access %5:array<array<i32, 4>, 4>, 0u, %4
+ ret %7
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func(%2:bool, %3:i32, %4:i32):i32 -> %b1 {
+ %b1 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block (%5:array<array<i32, 4>, 4>) { # body
+ %6:array<i32, 4> = access %5:array<array<i32, 4>, 4>, 0u
+ %7:ptr<function, array<i32, 4>, read_write> = var, %6
+ if %2 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ %8:ptr<function, i32, read_write> = access %7, %3
+ %9:i32 = load %8
+ ret %9
+ }
+ %b4 = block { # false
+ %10:ptr<function, i32, read_write> = access %7, %4
+ %11:i32 = load %10
+ ret %11
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest, MultipleAccessesToConstant_FromDifferentFunctions) {
+ auto* arr = b.Constant(mod.constant_values.Zero(ty.array<i32, 4>()));
+
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* func_a = b.Function("func_a", ty.i32());
+ func_a->SetParams({idx_a});
+ b.Append(func_a->Block(), [&] { //
+ b.Return(func_a, b.Access(ty.i32(), arr, idx_a));
+ });
+
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func_b = b.Function("func_b", ty.i32());
+ func_b->SetParams({idx_b});
+ b.Append(func_b->Block(), [&] { //
+ b.Return(func_b, b.Access(ty.i32(), arr, idx_b));
+ });
+
+ auto* idx_c = b.FunctionParam(ty.i32());
+ auto* func_c = b.Function("func_c", ty.i32());
+ func_c->SetParams({idx_c});
+ b.Append(func_c->Block(), [&] { //
+ b.Return(func_c, b.Access(ty.i32(), arr, idx_c));
+ });
+
+ auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+ %b1 = block {
+ %3:i32 = access array<i32, 4>(0i), %2
+ ret %3
+ }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+ %b2 = block {
+ %6:i32 = access array<i32, 4>(0i), %5
+ ret %6
+ }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+ %b3 = block {
+ %9:i32 = access array<i32, 4>(0i), %8
+ ret %9
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+ %b2 = block {
+ %4:ptr<private, i32, read_write> = access %1, %3
+ %5:i32 = load %4
+ ret %5
+ }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+ %b3 = block {
+ %8:ptr<private, i32, read_write> = access %1, %7
+ %9:i32 = load %8
+ ret %9
+ }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+ %b4 = block {
+ %12:ptr<private, i32, read_write> = access %1, %11
+ %13:i32 = load %12
+ ret %13
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_VarForDynamicIndexTest,
+ MultipleAccessesToConstant_FromDifferentFunctions_WithLeadingConstantIndex) {
+ auto* arr = b.Constant(mod.constant_values.Zero(ty.array(ty.array<i32, 4>(), 4)));
+
+ auto* idx_a = b.FunctionParam(ty.i32());
+ auto* func_a = b.Function("func_a", ty.i32());
+ func_a->SetParams({idx_a});
+ b.Append(func_a->Block(), [&] { //
+ b.Return(func_a, b.Access(ty.i32(), arr, 0_u, idx_a));
+ });
+
+ auto* idx_b = b.FunctionParam(ty.i32());
+ auto* func_b = b.Function("func_b", ty.i32());
+ func_b->SetParams({idx_b});
+ b.Append(func_b->Block(), [&] { //
+ b.Return(func_b, b.Access(ty.i32(), arr, 0_u, idx_b));
+ });
+
+ auto* idx_c = b.FunctionParam(ty.i32());
+ auto* func_c = b.Function("func_c", ty.i32());
+ func_c->SetParams({idx_c});
+ b.Append(func_c->Block(), [&] { //
+ b.Return(func_c, b.Access(ty.i32(), arr, 0_u, idx_c));
+ });
+
+ auto* src = R"(
+%func_a = func(%2:i32):i32 -> %b1 {
+ %b1 = block {
+ %3:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %2
+ ret %3
+ }
+}
+%func_b = func(%5:i32):i32 -> %b2 {
+ %b2 = block {
+ %6:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %5
+ ret %6
+ }
+}
+%func_c = func(%8:i32):i32 -> %b3 {
+ %b3 = block {
+ %9:i32 = access array<array<i32, 4>, 4>(array<i32, 4>(0i)), 0u, %8
+ ret %9
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%b1 = block { # root
+ %1:ptr<private, array<i32, 4>, read_write> = var, array<i32, 4>(0i)
+}
+
+%func_a = func(%3:i32):i32 -> %b2 {
+ %b2 = block {
+ %4:ptr<private, i32, read_write> = access %1, %3
+ %5:i32 = load %4
+ ret %5
+ }
+}
+%func_b = func(%7:i32):i32 -> %b3 {
+ %b3 = block {
+ %8:ptr<private, i32, read_write> = access %1, %7
+ %9:i32 = load %8
+ ret %9
+ }
+}
+%func_c = func(%11:i32):i32 -> %b4 {
+ %b4 = block {
+ %12:ptr<private, i32, read_write> = access %1, %11
+ %13:i32 = load %12
+ ret %13
+ }
+}
+)";
+
+ Run(VarForDynamicIndex);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::writer::raise
diff --git a/test/tint/bug/tint/2237.wgsl b/test/tint/bug/tint/2237.wgsl
new file mode 100644
index 0000000..31a667f
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl
@@ -0,0 +1,14 @@
+@group(0) @binding(0)
+var<storage, read_write> buffer : u32;
+
+const kArray = array(0u, 1u, 2u, 4u);
+
+fn foo() -> u32 {
+ return kArray[buffer];
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ let v = kArray[buffer];
+ buffer = v + foo();
+}
diff --git a/test/tint/bug/tint/2237.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2237.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..e7d539c
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.dxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer buffer : register(u0);
+
+uint foo() {
+ uint tint_symbol_2[4] = {0u, 1u, 2u, 4u};
+ return tint_symbol_2[buffer.Load(0u)];
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_3[4] = {0u, 1u, 2u, 4u};
+ uint v = tint_symbol_3[buffer.Load(0u)];
+ uint tint_symbol = v;
+ uint tint_symbol_1 = foo();
+ buffer.Store(0u, asuint((tint_symbol + tint_symbol_1)));
+ return;
+}
diff --git a/test/tint/bug/tint/2237.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2237.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..e7d539c
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.fxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer buffer : register(u0);
+
+uint foo() {
+ uint tint_symbol_2[4] = {0u, 1u, 2u, 4u};
+ return tint_symbol_2[buffer.Load(0u)];
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_3[4] = {0u, 1u, 2u, 4u};
+ uint v = tint_symbol_3[buffer.Load(0u)];
+ uint tint_symbol = v;
+ uint tint_symbol_1 = foo();
+ buffer.Store(0u, asuint((tint_symbol + tint_symbol_1)));
+ return;
+}
diff --git a/test/tint/bug/tint/2237.wgsl.expected.glsl b/test/tint/bug/tint/2237.wgsl.expected.glsl
new file mode 100644
index 0000000..c7d0d46
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ uint inner;
+} tint_symbol;
+
+uint foo() {
+ uint tint_symbol_4[4] = uint[4](0u, 1u, 2u, 4u);
+ return tint_symbol_4[tint_symbol.inner];
+}
+
+void tint_symbol_1() {
+ uint tint_symbol_5[4] = uint[4](0u, 1u, 2u, 4u);
+ uint v = tint_symbol_5[tint_symbol.inner];
+ uint tint_symbol_2 = v;
+ uint tint_symbol_3 = foo();
+ tint_symbol.inner = (tint_symbol_2 + tint_symbol_3);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1();
+ return;
+}
diff --git a/test/tint/bug/tint/2237.wgsl.expected.ir.spvasm b/test/tint/bug/tint/2237.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..fb6a04d
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.ir.spvasm
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 33
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %foo "foo"
+ OpName %main "main"
+ OpName %v "v"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %_arr_uint_uint_4 ArrayStride 4
+ %uint = OpTypeInt 32 0
+%tint_symbol_1 = OpTypeStruct %uint
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+%_ptr_Private__arr_uint_uint_4 = OpTypePointer Private %_arr_uint_uint_4
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %9 = OpConstantComposite %_arr_uint_uint_4 %uint_0 %uint_1 %uint_2 %uint_4
+ %5 = OpVariable %_ptr_Private__arr_uint_uint_4 Private %9
+ %14 = OpTypeFunction %uint
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_Private_uint = OpTypePointer Private %uint
+ %void = OpTypeVoid
+ %24 = OpTypeFunction %void
+ %foo = OpFunction %uint None %14
+ %15 = OpLabel
+ %16 = OpAccessChain %_ptr_StorageBuffer_uint %1 %uint_0
+ %18 = OpLoad %uint %16
+ %19 = OpAccessChain %_ptr_Private_uint %5 %18
+ %21 = OpLoad %uint %19
+ OpReturnValue %21
+ OpFunctionEnd
+ %main = OpFunction %void None %24
+ %25 = OpLabel
+ %26 = OpAccessChain %_ptr_StorageBuffer_uint %1 %uint_0
+ %27 = OpLoad %uint %26
+ %28 = OpAccessChain %_ptr_Private_uint %5 %27
+ %v = OpLoad %uint %28
+ %30 = OpFunctionCall %uint %foo
+ %31 = OpIAdd %uint %v %30
+ %32 = OpAccessChain %_ptr_StorageBuffer_uint %1 %uint_0
+ OpStore %32 %31
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2237.wgsl.expected.msl b/test/tint/bug/tint/2237.wgsl.expected.msl
new file mode 100644
index 0000000..b3a137b
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.msl
@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+uint foo(device uint* const tint_symbol_6) {
+ tint_array<uint, 4> const tint_symbol_4 = tint_array<uint, 4>{0u, 1u, 2u, 4u};
+ return tint_symbol_4[*(tint_symbol_6)];
+}
+
+kernel void tint_symbol_1(device uint* tint_symbol_7 [[buffer(0)]]) {
+ tint_array<uint, 4> const tint_symbol_5 = tint_array<uint, 4>{0u, 1u, 2u, 4u};
+ uint const v = tint_symbol_5[*(tint_symbol_7)];
+ uint const tint_symbol_2 = v;
+ uint const tint_symbol_3 = foo(tint_symbol_7);
+ *(tint_symbol_7) = (tint_symbol_2 + tint_symbol_3);
+ return;
+}
+
diff --git a/test/tint/bug/tint/2237.wgsl.expected.spvasm b/test/tint/bug/tint/2237.wgsl.expected.spvasm
new file mode 100644
index 0000000..49275d0
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %buffer_block "buffer_block"
+ OpMemberName %buffer_block 0 "inner"
+ OpName %buffer "buffer"
+ OpName %foo "foo"
+ OpName %var_for_index "var_for_index"
+ OpName %main "main"
+ OpName %var_for_index_1 "var_for_index_1"
+ OpDecorate %buffer_block Block
+ OpMemberDecorate %buffer_block 0 Offset 0
+ OpDecorate %buffer DescriptorSet 0
+ OpDecorate %buffer Binding 0
+ OpDecorate %_arr_uint_uint_4 ArrayStride 4
+ %uint = OpTypeInt 32 0
+%buffer_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_buffer_block = OpTypePointer StorageBuffer %buffer_block
+ %buffer = OpVariable %_ptr_StorageBuffer_buffer_block StorageBuffer
+ %5 = OpTypeFunction %uint
+ %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+ %10 = OpConstantNull %uint
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %13 = OpConstantComposite %_arr_uint_uint_4 %10 %uint_1 %uint_2 %uint_4
+%_ptr_Function__arr_uint_uint_4 = OpTypePointer Function %_arr_uint_uint_4
+ %16 = OpConstantNull %_arr_uint_uint_4
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %void = OpTypeVoid
+ %24 = OpTypeFunction %void
+ %foo = OpFunction %uint None %5
+ %7 = OpLabel
+%var_for_index = OpVariable %_ptr_Function__arr_uint_uint_4 Function %16
+ OpStore %var_for_index %13
+ %19 = OpAccessChain %_ptr_StorageBuffer_uint %buffer %uint_0
+ %20 = OpLoad %uint %19
+ %22 = OpAccessChain %_ptr_Function_uint %var_for_index %20
+ %23 = OpLoad %uint %22
+ OpReturnValue %23
+ OpFunctionEnd
+ %main = OpFunction %void None %24
+ %27 = OpLabel
+%var_for_index_1 = OpVariable %_ptr_Function__arr_uint_uint_4 Function %16
+ OpStore %var_for_index_1 %13
+ %29 = OpAccessChain %_ptr_StorageBuffer_uint %buffer %uint_0
+ %30 = OpLoad %uint %29
+ %31 = OpAccessChain %_ptr_Function_uint %var_for_index_1 %30
+ %32 = OpLoad %uint %31
+ %33 = OpFunctionCall %uint %foo
+ %34 = OpAccessChain %_ptr_StorageBuffer_uint %buffer %uint_0
+ %35 = OpIAdd %uint %32 %33
+ OpStore %34 %35
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2237.wgsl.expected.wgsl b/test/tint/bug/tint/2237.wgsl.expected.wgsl
new file mode 100644
index 0000000..92be5a6
--- /dev/null
+++ b/test/tint/bug/tint/2237.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+@group(0) @binding(0) var<storage, read_write> buffer : u32;
+
+const kArray = array(0u, 1u, 2u, 4u);
+
+fn foo() -> u32 {
+ return kArray[buffer];
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ let v = kArray[buffer];
+ buffer = (v + foo());
+}
