[hlsl] Fix decompose with reused access chain.
Currently the `Decompose*Access` transforms update the offset
information with each usage. This isn't correct as an access chain can
be re-used and subsequent usages need the original offset, not the
updated one. This Cl fixes both transforms to correctly calculate
offsets from re-used access chains.
Bug: 356631361
Change-Id: Ib8584163209941c3b7e696130ad360b9df6e5277
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/200917
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/hlsl/writer/access_test.cc b/src/tint/lang/hlsl/writer/access_test.cc
index f2d4120..194041c 100644
--- a/src/tint/lang/hlsl/writer/access_test.cc
+++ b/src/tint/lang/hlsl/writer/access_test.cc
@@ -2169,5 +2169,65 @@
)");
}
+TEST_F(HlslWriterTest, AccessChainReused) {
+ auto* sb = ty.Struct(mod.symbols.New("SB"), {
+ {mod.symbols.New("a"), ty.i32()},
+ {mod.symbols.New("b"), ty.vec3<f32>()},
+ });
+
+ auto* var = b.Var("v", storage, sb, core::Access::kReadWrite);
+ var->SetBindingPoint(0, 0);
+ b.ir.root_block->Append(var);
+
+ auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(func->Block(), [&] {
+ auto* x = b.Access(ty.ptr(storage, ty.vec3<f32>(), core::Access::kReadWrite), var, 1_u);
+ b.Let("b", b.LoadVectorElement(x, 1_u));
+ b.Let("c", b.LoadVectorElement(x, 2_u));
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
+ EXPECT_EQ(output_.hlsl, R"(
+RWByteAddressBuffer v : register(u0);
+void foo() {
+ float b = asfloat(v.Load(20u));
+ float c = asfloat(v.Load(24u));
+}
+
+)");
+}
+
+TEST_F(HlslWriterTest, UniformAccessChainReused) {
+ auto* sb = ty.Struct(mod.symbols.New("SB"), {
+ {mod.symbols.New("c"), ty.f32()},
+ {mod.symbols.New("d"), ty.vec3<f32>()},
+ });
+
+ auto* var = b.Var("v", uniform, sb, core::Access::kRead);
+ var->SetBindingPoint(0, 0);
+ b.ir.root_block->Append(var);
+
+ auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(func->Block(), [&] {
+ auto* x = b.Access(ty.ptr(uniform, ty.vec3<f32>(), core::Access::kRead), var, 1_u);
+ b.Let("b", b.LoadVectorElement(x, 1_u));
+ b.Let("c", b.LoadVectorElement(x, 2_u));
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
+ EXPECT_EQ(output_.hlsl, R"(
+cbuffer cbuffer_v : register(b0) {
+ uint4 v[2];
+};
+void foo() {
+ float b = asfloat(v[1u].y);
+ float c = asfloat(v[1u].z);
+}
+
+)");
+}
+
} // namespace
} // namespace tint::hlsl::writer
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc b/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
index 4f8b9f7..72198db 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_storage_access.cc
@@ -142,7 +142,7 @@
},
[&](core::ir::Access* a) {
OffsetData offset{};
- Access(a, var, a->Object()->Type(), &offset);
+ Access(a, var, a->Object()->Type(), offset);
},
[&](core::ir::Let* let) {
// The `let` is, essentially, an alias for the `var` as it's assigned
@@ -386,8 +386,6 @@
return;
}
- // TODO(dsinclair): This is missing arrays and matrices
-
tint::Switch(
from->Type(), //
[&](const core::type::Struct* s) {
@@ -703,9 +701,7 @@
void Access(core::ir::Access* a,
core::ir::Var* var,
const core::type::Type* obj,
- OffsetData* offset) {
- TINT_ASSERT(offset);
-
+ OffsetData offset) {
// Note, because we recurse through the `access` helper, the object passed in isn't
// necessarily the originating `var` object, but maybe a partially resolved access chain
// object.
@@ -717,17 +713,17 @@
tint::Switch(
obj, //
[&](const core::type::Vector* v) {
- b.InsertBefore(a,
- [&] { UpdateOffsetData(idx_value, v->type()->Size(), offset); });
+ b.InsertBefore(
+ a, [&] { UpdateOffsetData(idx_value, v->type()->Size(), &offset); });
obj = v->type();
},
[&](const core::type::Matrix* m) {
- b.InsertBefore(a,
- [&] { UpdateOffsetData(idx_value, m->ColumnStride(), offset); });
+ b.InsertBefore(
+ a, [&] { UpdateOffsetData(idx_value, m->ColumnStride(), &offset); });
obj = m->ColumnType();
},
[&](const core::type::Array* ary) {
- b.InsertBefore(a, [&] { UpdateOffsetData(idx_value, ary->Stride(), offset); });
+ b.InsertBefore(a, [&] { UpdateOffsetData(idx_value, ary->Stride(), &offset); });
obj = ary->ElemType();
},
[&](const core::type::Struct* s) {
@@ -738,7 +734,7 @@
uint32_t idx = cnst->Value()->ValueAs<uint32_t>();
auto* mem = s->Members()[idx];
- offset->byte_offset += mem->Offset();
+ offset.byte_offset += mem->Offset();
obj = mem->Type();
},
TINT_ICE_ON_NO_MATCH);
@@ -770,25 +766,28 @@
[&](core::ir::LoadVectorElement* lve) {
a->Result(0)->RemoveUsage(usage);
+ OffsetData load_offset = offset;
b.InsertBefore(lve, [&] {
- UpdateOffsetData(lve->Index(), obj->DeepestElement()->Size(), offset);
+ UpdateOffsetData(lve->Index(), obj->DeepestElement()->Size(), &load_offset);
});
- Load(lve, var, *offset);
+ Load(lve, var, load_offset);
},
[&](core::ir::Load* ld) {
a->Result(0)->RemoveUsage(usage);
- Load(ld, var, *offset);
+ Load(ld, var, offset);
},
[&](core::ir::StoreVectorElement* sve) {
a->Result(0)->RemoveUsage(usage);
+ OffsetData store_offset = offset;
b.InsertBefore(sve, [&] {
- UpdateOffsetData(sve->Index(), obj->DeepestElement()->Size(), offset);
+ UpdateOffsetData(sve->Index(), obj->DeepestElement()->Size(),
+ &store_offset);
});
- Store(sve, var, sve->Value(), *offset);
+ Store(sve, var, sve->Value(), store_offset);
},
- [&](core::ir::Store* store) { Store(store, var, store->From(), *offset); },
+ [&](core::ir::Store* store) { Store(store, var, store->From(), offset); },
[&](core::ir::CoreBuiltinCall* call) {
switch (call->Func()) {
case core::BuiltinFn::kArrayLength:
@@ -796,40 +795,40 @@
// access chain _must_ have resolved to the runtime array member of the
// structure. So, we _must_ have set `obj` to the array member which is
// a runtime array.
- ArrayLength(var, call, obj, offset->byte_offset);
+ ArrayLength(var, call, obj, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicAnd:
- AtomicAnd(var, call, offset->byte_offset);
+ AtomicAnd(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicOr:
- AtomicOr(var, call, offset->byte_offset);
+ AtomicOr(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicXor:
- AtomicXor(var, call, offset->byte_offset);
+ AtomicXor(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicMin:
- AtomicMin(var, call, offset->byte_offset);
+ AtomicMin(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicMax:
- AtomicMax(var, call, offset->byte_offset);
+ AtomicMax(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicAdd:
- AtomicAdd(var, call, offset->byte_offset);
+ AtomicAdd(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicSub:
- AtomicSub(var, call, offset->byte_offset);
+ AtomicSub(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicExchange:
- AtomicExchange(var, call, offset->byte_offset);
+ AtomicExchange(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicCompareExchangeWeak:
- AtomicCompareExchangeWeak(var, call, offset->byte_offset);
+ AtomicCompareExchangeWeak(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicStore:
- AtomicStore(var, call, offset->byte_offset);
+ AtomicStore(var, call, offset.byte_offset);
break;
case core::BuiltinFn::kAtomicLoad:
- AtomicLoad(var, call, offset->byte_offset);
+ AtomicLoad(var, call, offset.byte_offset);
break;
default:
TINT_UNREACHABLE() << call->Func();
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
index c868549..036562b 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
@@ -3579,5 +3579,73 @@
EXPECT_EQ(expect, str());
}
+TEST_F(HlslWriterDecomposeStorageAccessTest, AccessChainReused) {
+ auto* sb = ty.Struct(mod.symbols.New("SB"), {
+ {mod.symbols.New("a"), ty.i32()},
+ {mod.symbols.New("b"), ty.vec3<f32>()},
+ });
+
+ auto* var = b.Var("v", storage, sb, core::Access::kReadWrite);
+ var->SetBindingPoint(0, 0);
+ b.ir.root_block->Append(var);
+
+ auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(func->Block(), [&] {
+ auto* x = b.Access(ty.ptr(storage, ty.vec3<f32>(), core::Access::kReadWrite), var, 1_u);
+ b.Let("b", b.LoadVectorElement(x, 1_u));
+ b.Let("c", b.LoadVectorElement(x, 2_u));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+SB = struct @align(16) {
+ a:i32 @offset(0)
+ b:vec3<f32> @offset(16)
+}
+
+$B1: { # root
+ %v:ptr<storage, SB, read_write> = var @binding_point(0, 0)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:ptr<storage, vec3<f32>, read_write> = access %v, 1u
+ %4:f32 = load_vector_element %3, 1u
+ %b:f32 = let %4
+ %6:f32 = load_vector_element %3, 2u
+ %c:f32 = let %6
+ ret
+ }
+}
+)";
+ ASSERT_EQ(src, str());
+
+ auto* expect = R"(
+SB = struct @align(16) {
+ a:i32 @offset(0)
+ b:vec3<f32> @offset(16)
+}
+
+$B1: { # root
+ %v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:u32 = %v.Load 20u
+ %4:f32 = bitcast %3
+ %b:f32 = let %4
+ %6:u32 = %v.Load 24u
+ %7:f32 = bitcast %6
+ %c:f32 = let %7
+ ret
+ }
+}
+)";
+
+ Run(DecomposeStorageAccess);
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::hlsl::writer::raise
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_uniform_access.cc b/src/tint/lang/hlsl/writer/raise/decompose_uniform_access.cc
index 9fb8ff5..741d476 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_uniform_access.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_uniform_access.cc
@@ -173,7 +173,7 @@
void Access(core::ir::Access* a,
core::ir::Var* var,
const core::type::Type* obj_ty,
- OffsetData& offset) {
+ OffsetData offset) {
// Note, because we recurse through the `access` helper, the object passed in isn't
// necessarily the originating `var` object, but maybe a partially resolved access chain
// object.
@@ -259,7 +259,7 @@
a->Destroy();
}
- void Load(core::ir::Load* ld, core::ir::Var* var, OffsetData& offset) {
+ void Load(core::ir::Load* ld, core::ir::Var* var, OffsetData offset) {
b.InsertBefore(ld, [&] {
auto* byte_idx = OffsetToValue(offset);
auto* result = MakeLoad(ld, var, ld->Result(0)->Type(), byte_idx);
@@ -270,7 +270,7 @@
void LoadVectorElement(core::ir::LoadVectorElement* lve,
core::ir::Var* var,
- OffsetData& offset) {
+ OffsetData offset) {
b.InsertBefore(lve, [&] {
// Add the byte count from the start of the vector to the requested element to the
// current offset calculation
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_uniform_access_test.cc b/src/tint/lang/hlsl/writer/raise/decompose_uniform_access_test.cc
index 86b61e8..d2be685 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_uniform_access_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_uniform_access_test.cc
@@ -1050,5 +1050,75 @@
EXPECT_EQ(expect, str());
}
+TEST_F(HlslWriterDecomposeUniformAccessTest, UniformAccessChainReused) {
+ auto* sb = ty.Struct(mod.symbols.New("SB"), {
+ {mod.symbols.New("c"), ty.f32()},
+ {mod.symbols.New("d"), ty.vec3<f32>()},
+ });
+
+ auto* var = b.Var("v", uniform, sb, core::Access::kRead);
+ var->SetBindingPoint(0, 0);
+ b.ir.root_block->Append(var);
+
+ auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(func->Block(), [&] {
+ auto* x = b.Access(ty.ptr(uniform, ty.vec3<f32>(), core::Access::kRead), var, 1_u);
+ b.Let("b", b.LoadVectorElement(x, 1_u));
+ b.Let("c", b.LoadVectorElement(x, 2_u));
+ b.Return(func);
+ });
+
+ auto* src = R"(
+SB = struct @align(16) {
+ c:f32 @offset(0)
+ d:vec3<f32> @offset(16)
+}
+
+$B1: { # root
+ %v:ptr<uniform, SB, read> = var @binding_point(0, 0)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:ptr<uniform, vec3<f32>, read> = access %v, 1u
+ %4:f32 = load_vector_element %3, 1u
+ %b:f32 = let %4
+ %6:f32 = load_vector_element %3, 2u
+ %c:f32 = let %6
+ ret
+ }
+}
+)";
+ ASSERT_EQ(src, str());
+
+ auto* expect = R"(
+SB = struct @align(16) {
+ c:f32 @offset(0)
+ d:vec3<f32> @offset(16)
+}
+
+$B1: { # root
+ %v:ptr<uniform, array<vec4<u32>, 2>, read> = var @binding_point(0, 0)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:ptr<uniform, vec4<u32>, read> = access %v, 1u
+ %4:u32 = load_vector_element %3, 1u
+ %5:f32 = bitcast %4
+ %b:f32 = let %5
+ %7:ptr<uniform, vec4<u32>, read> = access %v, 1u
+ %8:u32 = load_vector_element %7, 2u
+ %9:f32 = bitcast %8
+ %c:f32 = let %9
+ ret
+ }
+}
+)";
+
+ Run(DecomposeUniformAccess);
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::hlsl::writer::raise
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index ab6bb80..2409d0b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -43,7 +43,7 @@
uint v_8 = (16u * uint(i()));
uint v_9 = (8u * uint(i()));
float2x2 v_10[4] = v_4(0u);
- float2x2 l_a_i = v((v_8 + v_9));
+ float2x2 l_a_i = v(v_8);
uint4 v_11 = a[((v_8 + v_9) / 16u)];
float2 l_a_i_i = asfloat(((((((v_8 + v_9) % 16u) / 4u) == 2u)) ? (v_11.zw) : (v_11.xy)));
float2x2 l_a[4] = v_10;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index d94780f..5e09c8b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -35,7 +35,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x2 v_8[4] = v_4(0u);
- float2x2 l_a_i = v(40u);
+ float2x2 l_a_i = v(32u);
float2 l_a_i_i = asfloat(a[2u].zw);
float2x2 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[2u].z) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index d94780f..5e09c8b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -35,7 +35,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x2 v_8[4] = v_4(0u);
- float2x2 l_a_i = v(40u);
+ float2x2 l_a_i = v(32u);
float2 l_a_i_i = asfloat(a[2u].zw);
float2x2 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[2u].z) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 77586a1..f6db147 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -53,7 +53,7 @@
uint v_10 = (16u * uint(i()));
uint v_11 = (8u * uint(i()));
matrix<float16_t, 2, 3> v_12[4] = v_6(0u);
- matrix<float16_t, 2, 3> l_a_i = v_4((v_10 + v_11));
+ matrix<float16_t, 2, 3> l_a_i = v_4(v_10);
vector<float16_t, 3> l_a_i_i = tint_bitcast_to_f16(a[((v_10 + v_11) / 16u)]).xyz;
uint v_13 = a[((v_10 + v_11) / 16u)][(((v_10 + v_11) % 16u) / 4u)];
matrix<float16_t, 2, 3> l_a[4] = v_12;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 5dc4645..3e44f36 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -45,7 +45,7 @@
[numthreads(1, 1, 1)]
void f() {
matrix<float16_t, 2, 3> v_10[4] = v_6(0u);
- matrix<float16_t, 2, 3> l_a_i = v_4(40u);
+ matrix<float16_t, 2, 3> l_a_i = v_4(32u);
vector<float16_t, 3> l_a_i_i = tint_bitcast_to_f16(a[2u]).xyz;
matrix<float16_t, 2, 3> l_a[4] = v_10;
s.Store<float16_t>(0u, (((float16_t(f16tof32(a[2u].z)) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 86711fd..7db710b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -41,7 +41,7 @@
uint v_6 = (32u * uint(i()));
uint v_7 = (16u * uint(i()));
float2x3 v_8[4] = v_2(0u);
- float2x3 l_a_i = v((v_6 + v_7));
+ float2x3 l_a_i = v(v_6);
float3 l_a_i_i = asfloat(a[((v_6 + v_7) / 16u)].xyz);
float2x3 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[((v_6 + v_7) / 16u)][(((v_6 + v_7) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 86711fd..7db710b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -41,7 +41,7 @@
uint v_6 = (32u * uint(i()));
uint v_7 = (16u * uint(i()));
float2x3 v_8[4] = v_2(0u);
- float2x3 l_a_i = v((v_6 + v_7));
+ float2x3 l_a_i = v(v_6);
float3 l_a_i_i = asfloat(a[((v_6 + v_7) / 16u)].xyz);
float2x3 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[((v_6 + v_7) / 16u)][(((v_6 + v_7) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index e8fa3b9..873afd1 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -33,7 +33,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x3 v_6[4] = v_2(0u);
- float2x3 l_a_i = v(80u);
+ float2x3 l_a_i = v(64u);
float3 l_a_i_i = asfloat(a[5u].xyz);
float2x3 l_a[4] = v_6;
s.Store(0u, asuint((((asfloat(a[5u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index e8fa3b9..873afd1 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -33,7 +33,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x3 v_6[4] = v_2(0u);
- float2x3 l_a_i = v(80u);
+ float2x3 l_a_i = v(64u);
float3 l_a_i_i = asfloat(a[5u].xyz);
float2x3 l_a[4] = v_6;
s.Store(0u, asuint((((asfloat(a[5u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 753c2d2..5677e6e 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -41,7 +41,7 @@
uint v_6 = (32u * uint(i()));
uint v_7 = (16u * uint(i()));
float2x4 v_8[4] = v_2(0u);
- float2x4 l_a_i = v((v_6 + v_7));
+ float2x4 l_a_i = v(v_6);
float4 l_a_i_i = asfloat(a[((v_6 + v_7) / 16u)]);
float2x4 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[((v_6 + v_7) / 16u)][(((v_6 + v_7) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index dbc3ee3..e4e0bd4 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -33,7 +33,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x4 v_6[4] = v_2(0u);
- float2x4 l_a_i = v(80u);
+ float2x4 l_a_i = v(64u);
float4 l_a_i_i = asfloat(a[5u]);
float2x4 l_a[4] = v_6;
s.Store(0u, asuint((((asfloat(a[5u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index dbc3ee3..e4e0bd4 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -33,7 +33,7 @@
[numthreads(1, 1, 1)]
void f() {
float2x4 v_6[4] = v_2(0u);
- float2x4 l_a_i = v(80u);
+ float2x4 l_a_i = v(64u);
float4 l_a_i_i = asfloat(a[5u]);
float2x4 l_a[4] = v_6;
s.Store(0u, asuint((((asfloat(a[5u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 23e3206..292e8a2 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -42,7 +42,7 @@
uint v_7 = (48u * uint(i()));
uint v_8 = (16u * uint(i()));
float3x3 v_9[4] = v_3(0u);
- float3x3 l_a_i = v((v_7 + v_8));
+ float3x3 l_a_i = v(v_7);
float3 l_a_i_i = asfloat(a[((v_7 + v_8) / 16u)].xyz);
float3x3 l_a[4] = v_9;
s.Store(0u, asuint((((asfloat(a[((v_7 + v_8) / 16u)][(((v_7 + v_8) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 9a00d43..063ff2e 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -34,7 +34,7 @@
[numthreads(1, 1, 1)]
void f() {
float3x3 v_7[4] = v_3(0u);
- float3x3 l_a_i = v(112u);
+ float3x3 l_a_i = v(96u);
float3 l_a_i_i = asfloat(a[7u].xyz);
float3x3 l_a[4] = v_7;
s.Store(0u, asuint((((asfloat(a[7u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 30f28c9..f3abf1b 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -42,7 +42,7 @@
uint v_7 = (48u * uint(i()));
uint v_8 = (16u * uint(i()));
float3x4 v_9[4] = v_3(0u);
- float3x4 l_a_i = v((v_7 + v_8));
+ float3x4 l_a_i = v(v_7);
float4 l_a_i_i = asfloat(a[((v_7 + v_8) / 16u)]);
float3x4 l_a[4] = v_9;
s.Store(0u, asuint((((asfloat(a[((v_7 + v_8) / 16u)][(((v_7 + v_8) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 30f28c9..f3abf1b 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -42,7 +42,7 @@
uint v_7 = (48u * uint(i()));
uint v_8 = (16u * uint(i()));
float3x4 v_9[4] = v_3(0u);
- float3x4 l_a_i = v((v_7 + v_8));
+ float3x4 l_a_i = v(v_7);
float4 l_a_i_i = asfloat(a[((v_7 + v_8) / 16u)]);
float3x4 l_a[4] = v_9;
s.Store(0u, asuint((((asfloat(a[((v_7 + v_8) / 16u)][(((v_7 + v_8) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 16725a1..bc1a892 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -34,7 +34,7 @@
[numthreads(1, 1, 1)]
void f() {
float3x4 v_7[4] = v_3(0u);
- float3x4 l_a_i = v(112u);
+ float3x4 l_a_i = v(96u);
float4 l_a_i_i = asfloat(a[7u]);
float3x4 l_a[4] = v_7;
s.Store(0u, asuint((((asfloat(a[7u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 60f1c9a..bde35e1 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -55,7 +55,7 @@
uint v_14 = (16u * uint(i()));
uint v_15 = (4u * uint(i()));
matrix<float16_t, 4, 2> v_16[4] = v_10(0u);
- matrix<float16_t, 4, 2> l_a_i = v_2((v_14 + v_15));
+ matrix<float16_t, 4, 2> l_a_i = v_2(v_14);
uint4 v_17 = a[((v_14 + v_15) / 16u)];
vector<float16_t, 2> l_a_i_i = tint_bitcast_to_f16(((((((v_14 + v_15) % 16u) / 4u) == 2u)) ? (v_17.z) : (v_17.x)));
uint v_18 = a[((v_14 + v_15) / 16u)][(((v_14 + v_15) % 16u) / 4u)];
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index a93e556..e3da59c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -47,7 +47,7 @@
uint v_12 = (32u * uint(i()));
uint v_13 = (8u * uint(i()));
float4x2 v_14[4] = v_8(0u);
- float4x2 l_a_i = v((v_12 + v_13));
+ float4x2 l_a_i = v(v_12);
uint4 v_15 = a[((v_12 + v_13) / 16u)];
float2 l_a_i_i = asfloat(((((((v_12 + v_13) % 16u) / 4u) == 2u)) ? (v_15.zw) : (v_15.xy)));
float4x2 l_a[4] = v_14;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index a576fdb..4d68b3a 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -39,7 +39,7 @@
[numthreads(1, 1, 1)]
void f() {
float4x2 v_12[4] = v_8(0u);
- float4x2 l_a_i = v(72u);
+ float4x2 l_a_i = v(64u);
float2 l_a_i_i = asfloat(a[4u].zw);
float4x2 l_a[4] = v_12;
s.Store(0u, asuint((((asfloat(a[4u].z) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 23eab21..ac4afe2 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -55,7 +55,7 @@
uint v_12 = (32u * uint(i()));
uint v_13 = (8u * uint(i()));
matrix<float16_t, 4, 3> v_14[4] = v_8(0u);
- matrix<float16_t, 4, 3> l_a_i = v_4((v_12 + v_13));
+ matrix<float16_t, 4, 3> l_a_i = v_4(v_12);
vector<float16_t, 3> l_a_i_i = tint_bitcast_to_f16(a[((v_12 + v_13) / 16u)]).xyz;
uint v_15 = a[((v_12 + v_13) / 16u)][(((v_12 + v_13) % 16u) / 4u)];
matrix<float16_t, 4, 3> l_a[4] = v_14;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 52a24b3..f232808 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -47,7 +47,7 @@
[numthreads(1, 1, 1)]
void f() {
matrix<float16_t, 4, 3> v_12[4] = v_8(0u);
- matrix<float16_t, 4, 3> l_a_i = v_4(72u);
+ matrix<float16_t, 4, 3> l_a_i = v_4(64u);
vector<float16_t, 3> l_a_i_i = tint_bitcast_to_f16(a[4u]).xyz;
matrix<float16_t, 4, 3> l_a[4] = v_12;
s.Store<float16_t>(0u, (((float16_t(f16tof32(a[4u].z)) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 8f61ff5..0a15e6c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -43,7 +43,7 @@
uint v_8 = (64u * uint(i()));
uint v_9 = (16u * uint(i()));
float4x3 v_10[4] = v_4(0u);
- float4x3 l_a_i = v((v_8 + v_9));
+ float4x3 l_a_i = v(v_8);
float3 l_a_i_i = asfloat(a[((v_8 + v_9) / 16u)].xyz);
float4x3 l_a[4] = v_10;
s.Store(0u, asuint((((asfloat(a[((v_8 + v_9) / 16u)][(((v_8 + v_9) % 16u) / 4u)]) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 2db9988..959612c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -35,7 +35,7 @@
[numthreads(1, 1, 1)]
void f() {
float4x3 v_8[4] = v_4(0u);
- float4x3 l_a_i = v(144u);
+ float4x3 l_a_i = v(128u);
float3 l_a_i_i = asfloat(a[9u].xyz);
float4x3 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[9u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 2db9988..959612c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -35,7 +35,7 @@
[numthreads(1, 1, 1)]
void f() {
float4x3 v_8[4] = v_4(0u);
- float4x3 l_a_i = v(144u);
+ float4x3 l_a_i = v(128u);
float3 l_a_i_i = asfloat(a[9u].xyz);
float4x3 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[9u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 11a4265..047ba3e 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -47,7 +47,7 @@
[numthreads(1, 1, 1)]
void f() {
matrix<float16_t, 4, 4> v_12[4] = v_8(0u);
- matrix<float16_t, 4, 4> l_a_i = v_4(72u);
+ matrix<float16_t, 4, 4> l_a_i = v_4(64u);
vector<float16_t, 4> l_a_i_i = tint_bitcast_to_f16(a[4u]);
matrix<float16_t, 4, 4> l_a[4] = v_12;
s.Store<float16_t>(0u, (((float16_t(f16tof32(a[4u].z)) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u]));
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index e50d125..908add0 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -35,7 +35,7 @@
[numthreads(1, 1, 1)]
void f() {
float4x4 v_8[4] = v_4(0u);
- float4x4 l_a_i = v(144u);
+ float4x4 l_a_i = v(128u);
float4 l_a_i_i = asfloat(a[9u]);
float4x4 l_a[4] = v_8;
s.Store(0u, asuint((((asfloat(a[9u].x) + l_a[0][0][0u]) + l_a_i[0][0u]) + l_a_i_i[0u])));
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index eb004f4..38cfc8d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -85,9 +85,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_16(0u);
- Outer l_a_3 = v_13(900u);
+ Outer l_a_3 = v_13(768u);
Inner l_a_3_a[4] = v_8(768u);
- Inner l_a_3_a_2 = v_6(900u);
+ Inner l_a_3_a_2 = v_6(896u);
matrix<float16_t, 2, 2> l_a_3_a_2_m = v_2(896u);
vector<float16_t, 2> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u].x);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].y));
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 711901c..a18c41f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -77,9 +77,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_14(0u);
- Outer l_a_3 = v_11(904u);
+ Outer l_a_3 = v_11(768u);
Inner l_a_3_a[4] = v_6(768u);
- Inner l_a_3_a_2 = v_4(904u);
+ Inner l_a_3_a_2 = v_4(896u);
float2x2 l_a_3_a_2_m = v(896u);
float2 l_a_3_a_2_m_1 = asfloat(a[56u].zw);
float l_a_3_a_2_m_1_0 = asfloat(a[56u].z);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 3a4dabd..a18c41f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -77,10 +77,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_14(0u);
- Outer l_a_3 = v_11(904u);
- Inner l_a_3_a[4] = v_6(904u);
- Inner l_a_3_a_2 = v_4(904u);
- float2x2 l_a_3_a_2_m = v(904u);
+ Outer l_a_3 = v_11(768u);
+ Inner l_a_3_a[4] = v_6(768u);
+ Inner l_a_3_a_2 = v_4(896u);
+ float2x2 l_a_3_a_2_m = v(896u);
float2 l_a_3_a_2_m_1 = asfloat(a[56u].zw);
float l_a_3_a_2_m_1_0 = asfloat(a[56u].z);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index f8d0da8..7809a22 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -88,9 +88,9 @@
void f() {
Outer l_a[4] = v_16(0u);
Outer l_a_3 = v_13(768u);
- Inner l_a_3_a[4] = v_8(904u);
- Inner l_a_3_a_2 = v_6(904u);
- matrix<float16_t, 2, 3> l_a_3_a_2_m = v_4(904u);
+ Inner l_a_3_a[4] = v_8(768u);
+ Inner l_a_3_a_2 = v_6(896u);
+ matrix<float16_t, 2, 3> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 3> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]).xyz;
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 43359c6..e736df3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -75,9 +75,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_12(0u);
- Outer l_a_3 = v_9(912u);
- Inner l_a_3_a[4] = v_4(912u);
- Inner l_a_3_a_2 = v_2(912u);
+ Outer l_a_3 = v_9(768u);
+ Inner l_a_3_a[4] = v_4(768u);
+ Inner l_a_3_a_2 = v_2(896u);
float2x3 l_a_3_a_2_m = v(896u);
float3 l_a_3_a_2_m_1 = asfloat(a[57u].xyz);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index b5a8c79..e736df3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -75,10 +75,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_12(0u);
- Outer l_a_3 = v_9(912u);
+ Outer l_a_3 = v_9(768u);
Inner l_a_3_a[4] = v_4(768u);
Inner l_a_3_a_2 = v_2(896u);
- float2x3 l_a_3_a_2_m = v(912u);
+ float2x3 l_a_3_a_2_m = v(896u);
float3 l_a_3_a_2_m_1 = asfloat(a[57u].xyz);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 3dd0457..c304e25 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -87,9 +87,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_16(0u);
- Outer l_a_3 = v_13(904u);
- Inner l_a_3_a[4] = v_8(904u);
- Inner l_a_3_a_2 = v_6(904u);
+ Outer l_a_3 = v_13(768u);
+ Inner l_a_3_a[4] = v_8(768u);
+ Inner l_a_3_a_2 = v_6(896u);
matrix<float16_t, 2, 4> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 4> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index c688158..422aacc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -75,9 +75,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_12(0u);
- Outer l_a_3 = v_9(912u);
- Inner l_a_3_a[4] = v_4(912u);
- Inner l_a_3_a_2 = v_2(912u);
+ Outer l_a_3 = v_9(768u);
+ Inner l_a_3_a[4] = v_4(768u);
+ Inner l_a_3_a_2 = v_2(896u);
float2x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 563b222..422aacc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -75,10 +75,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_12(0u);
- Outer l_a_3 = v_9(912u);
- Inner l_a_3_a[4] = v_4(912u);
- Inner l_a_3_a_2 = v_2(912u);
- float2x4 l_a_3_a_2_m = v(912u);
+ Outer l_a_3 = v_9(768u);
+ Inner l_a_3_a[4] = v_4(768u);
+ Inner l_a_3_a_2 = v_2(896u);
+ float2x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 8f7b99b..a1e3996 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -88,9 +88,9 @@
void f() {
Outer l_a[4] = v_18(0u);
Outer l_a_3 = v_15(768u);
- Inner l_a_3_a[4] = v_10(900u);
- Inner l_a_3_a_2 = v_8(900u);
- matrix<float16_t, 3, 2> l_a_3_a_2_m = v_2(900u);
+ Inner l_a_3_a[4] = v_10(768u);
+ Inner l_a_3_a_2 = v_8(896u);
+ matrix<float16_t, 3, 2> l_a_3_a_2_m = v_2(896u);
vector<float16_t, 2> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u].x);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].y));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index d097d16..1f070fa 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -79,8 +79,8 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_16(0u);
- Outer l_a_3 = v_13(904u);
- Inner l_a_3_a[4] = v_8(904u);
+ Outer l_a_3 = v_13(768u);
+ Inner l_a_3_a[4] = v_8(768u);
Inner l_a_3_a_2 = v_6(896u);
float3x2 l_a_3_a_2_m = v(896u);
float2 l_a_3_a_2_m_1 = asfloat(a[56u].zw);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 817b99c..e66934f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -90,7 +90,7 @@
Outer l_a[4] = v_17(0u);
Outer l_a_3 = v_14(768u);
Inner l_a_3_a[4] = v_9(768u);
- Inner l_a_3_a_2 = v_7(904u);
+ Inner l_a_3_a_2 = v_7(896u);
matrix<float16_t, 3, 3> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 3> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]).xyz;
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index c6f7686..6460b52 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -76,7 +76,7 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_13(0u);
- Outer l_a_3 = v_10(912u);
+ Outer l_a_3 = v_10(768u);
Inner l_a_3_a[4] = v_5(768u);
Inner l_a_3_a_2 = v_3(896u);
float3x3 l_a_3_a_2_m = v(896u);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 00b58c2..6460b52 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -77,9 +77,9 @@
void f() {
Outer l_a[4] = v_13(0u);
Outer l_a_3 = v_10(768u);
- Inner l_a_3_a[4] = v_5(912u);
+ Inner l_a_3_a[4] = v_5(768u);
Inner l_a_3_a_2 = v_3(896u);
- float3x3 l_a_3_a_2_m = v(912u);
+ float3x3 l_a_3_a_2_m = v(896u);
float3 l_a_3_a_2_m_1 = asfloat(a[57u].xyz);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 1779345..c5013aa 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -88,10 +88,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_17(0u);
- Outer l_a_3 = v_14(904u);
- Inner l_a_3_a[4] = v_9(904u);
+ Outer l_a_3 = v_14(768u);
+ Inner l_a_3_a[4] = v_9(768u);
Inner l_a_3_a_2 = v_7(896u);
- matrix<float16_t, 3, 4> l_a_3_a_2_m = v_4(904u);
+ matrix<float16_t, 3, 4> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 4> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index e0c8207..3a92bc7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -76,9 +76,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_13(0u);
- Outer l_a_3 = v_10(912u);
+ Outer l_a_3 = v_10(768u);
Inner l_a_3_a[4] = v_5(768u);
- Inner l_a_3_a_2 = v_3(912u);
+ Inner l_a_3_a_2 = v_3(896u);
float3x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index e0c8207..3a92bc7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -76,9 +76,9 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_13(0u);
- Outer l_a_3 = v_10(912u);
+ Outer l_a_3 = v_10(768u);
Inner l_a_3_a[4] = v_5(768u);
- Inner l_a_3_a_2 = v_3(912u);
+ Inner l_a_3_a_2 = v_3(896u);
float3x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index b7bf697..ebac656 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -90,9 +90,9 @@
void f() {
Outer l_a[4] = v_20(0u);
Outer l_a_3 = v_17(768u);
- Inner l_a_3_a[4] = v_12(900u);
+ Inner l_a_3_a[4] = v_12(768u);
Inner l_a_3_a_2 = v_10(896u);
- matrix<float16_t, 4, 2> l_a_3_a_2_m = v_2(900u);
+ matrix<float16_t, 4, 2> l_a_3_a_2_m = v_2(896u);
vector<float16_t, 2> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u].x);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].y));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 157a36b..2ab1fbd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -82,7 +82,7 @@
void f() {
Outer l_a[4] = v_18(0u);
Outer l_a_3 = v_15(768u);
- Inner l_a_3_a[4] = v_10(904u);
+ Inner l_a_3_a[4] = v_10(768u);
Inner l_a_3_a_2 = v_8(896u);
float4x2 l_a_3_a_2_m = v(896u);
float2 l_a_3_a_2_m_1 = asfloat(a[56u].zw);
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 9ea1317..2ab1fbd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -81,10 +81,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_18(0u);
- Outer l_a_3 = v_15(904u);
- Inner l_a_3_a[4] = v_10(904u);
+ Outer l_a_3 = v_15(768u);
+ Inner l_a_3_a[4] = v_10(768u);
Inner l_a_3_a_2 = v_8(896u);
- float4x2 l_a_3_a_2_m = v(904u);
+ float4x2 l_a_3_a_2_m = v(896u);
float2 l_a_3_a_2_m_1 = asfloat(a[56u].zw);
float l_a_3_a_2_m_1_0 = asfloat(a[56u].z);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 87c8762..01abb40 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -89,10 +89,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_18(0u);
- Outer l_a_3 = v_15(904u);
- Inner l_a_3_a[4] = v_10(904u);
- Inner l_a_3_a_2 = v_8(904u);
- matrix<float16_t, 4, 3> l_a_3_a_2_m = v_4(904u);
+ Outer l_a_3 = v_15(768u);
+ Inner l_a_3_a[4] = v_10(768u);
+ Inner l_a_3_a_2 = v_8(896u);
+ matrix<float16_t, 4, 3> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 3> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]).xyz;
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 3edba00..c2ba7a9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -78,9 +78,9 @@
void f() {
Outer l_a[4] = v_14(0u);
Outer l_a_3 = v_11(768u);
- Inner l_a_3_a[4] = v_6(912u);
+ Inner l_a_3_a[4] = v_6(768u);
Inner l_a_3_a_2 = v_4(896u);
- float4x3 l_a_3_a_2_m = v(912u);
+ float4x3 l_a_3_a_2_m = v(896u);
float3 l_a_3_a_2_m_1 = asfloat(a[57u].xyz);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 60577b9..c2ba7a9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -77,10 +77,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_14(0u);
- Outer l_a_3 = v_11(912u);
- Inner l_a_3_a[4] = v_6(912u);
- Inner l_a_3_a_2 = v_4(912u);
- float4x3 l_a_3_a_2_m = v(912u);
+ Outer l_a_3 = v_11(768u);
+ Inner l_a_3_a[4] = v_6(768u);
+ Inner l_a_3_a_2 = v_4(896u);
+ float4x3 l_a_3_a_2_m = v(896u);
float3 l_a_3_a_2_m_1 = asfloat(a[57u].xyz);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 113ba8c..8b3fadb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -89,10 +89,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_18(0u);
- Outer l_a_3 = v_15(904u);
- Inner l_a_3_a[4] = v_10(904u);
- Inner l_a_3_a_2 = v_8(904u);
- matrix<float16_t, 4, 4> l_a_3_a_2_m = v_4(904u);
+ Outer l_a_3 = v_15(768u);
+ Inner l_a_3_a[4] = v_10(768u);
+ Inner l_a_3_a_2 = v_8(896u);
+ matrix<float16_t, 4, 4> l_a_3_a_2_m = v_4(896u);
vector<float16_t, 4> l_a_3_a_2_m_1 = tint_bitcast_to_f16(a[56u]);
float16_t l_a_3_a_2_m_1_0 = float16_t(f16tof32(a[56u].z));
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
index 26278de..08d4c5d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.dxc.hlsl
@@ -77,10 +77,10 @@
[numthreads(1, 1, 1)]
void f() {
Outer l_a[4] = v_14(0u);
- Outer l_a_3 = v_11(912u);
- Inner l_a_3_a[4] = v_6(912u);
- Inner l_a_3_a_2 = v_4(912u);
- float4x4 l_a_3_a_2_m = v(912u);
+ Outer l_a_3 = v_11(768u);
+ Inner l_a_3_a[4] = v_6(768u);
+ Inner l_a_3_a_2 = v_4(896u);
+ float4x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
float l_a_3_a_2_m_1_0 = asfloat(a[57u].x);
}
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
index 891fd7b..08d4c5d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.fxc.hlsl
@@ -78,7 +78,7 @@
void f() {
Outer l_a[4] = v_14(0u);
Outer l_a_3 = v_11(768u);
- Inner l_a_3_a[4] = v_6(912u);
+ Inner l_a_3_a[4] = v_6(768u);
Inner l_a_3_a_2 = v_4(896u);
float4x4 l_a_3_a_2_m = v(896u);
float4 l_a_3_a_2_m_1 = asfloat(a[57u]);
diff --git a/test/tint/bug/chromium/344265982.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/chromium/344265982.wgsl.expected.ir.dxc.hlsl
index de5e6ea..e5f933b 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.ir.dxc.hlsl
@@ -8,8 +8,7 @@
} else {
break;
}
- uint v = (uint(i) * 4u);
- switch(asint(buffer.Load((0u + v)))) {
+ switch(asint(buffer.Load((0u + (uint(i) * 4u))))) {
case 1:
{
{
@@ -19,8 +18,8 @@
}
default:
{
- uint v_1 = ((0u + v) + (uint(i) * 4u));
- buffer.Store(v_1, asuint(2));
+ uint v = (0u + (uint(i) * 4u));
+ buffer.Store(v, asuint(2));
break;
}
}
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.dxc.hlsl b/test/tint/statements/decrement/complex.wgsl.expected.ir.dxc.hlsl
index 324a03d..f7a5474 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.dxc.hlsl
@@ -38,25 +38,23 @@
uint v_3 = (uint(v_1) * 64u);
uint v_4 = (uint(v_2) * 16u);
int v_5 = idx3();
- uint v_6 = (uint(v_5) * 4u);
- int v_7 = (asint(buffer.Load((((0u + v_3) + v_4) + v_6))) - 1);
- uint v_8 = ((((0u + v_3) + v_4) + v_6) + (uint(v_5) * 4u));
- buffer.Store(v_8, asuint(v_7));
+ int v_6 = (asint(buffer.Load((((0u + v_3) + v_4) + (uint(v_5) * 4u)))) - 1);
+ uint v_7 = (((0u + v_3) + v_4) + (uint(v_5) * 4u));
+ buffer.Store(v_7, asuint(v_6));
while(true) {
if ((v < 10u)) {
} else {
break;
}
{
- int v_9 = idx4();
- int v_10 = idx5();
- uint v_11 = (uint(v_9) * 64u);
- uint v_12 = (uint(v_10) * 16u);
- int v_13 = idx6();
- uint v_14 = (uint(v_13) * 4u);
- int v_15 = (asint(buffer.Load((((0u + v_11) + v_12) + v_14))) - 1);
- uint v_16 = ((((0u + v_11) + v_12) + v_14) + (uint(v_13) * 4u));
- buffer.Store(v_16, asuint(v_15));
+ int v_8 = idx4();
+ int v_9 = idx5();
+ uint v_10 = (uint(v_8) * 64u);
+ uint v_11 = (uint(v_9) * 16u);
+ int v_12 = idx6();
+ int v_13 = (asint(buffer.Load((((0u + v_10) + v_11) + (uint(v_12) * 4u)))) - 1);
+ uint v_14 = (((0u + v_10) + v_11) + (uint(v_12) * 4u));
+ buffer.Store(v_14, asuint(v_13));
}
continue;
}
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.fxc.hlsl b/test/tint/statements/decrement/complex.wgsl.expected.ir.fxc.hlsl
index ac360ca..f7a5474 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.fxc.hlsl
@@ -1,206 +1,67 @@
-SKIP: FAILED
-
-struct S {
- a : array<vec4<i32>, 4>,
-}
-
-@group(0) @binding(0) var<storage, read_write> buffer : array<S>;
-
-var<private> v : u32;
-
-fn idx1() -> i32 {
- v--;
+RWByteAddressBuffer buffer : register(u0);
+static uint v = 0u;
+int idx1() {
+ v = (v - 1u);
return 1;
}
-fn idx2() -> i32 {
- v--;
+int idx2() {
+ v = (v - 1u);
return 2;
}
-fn idx3() -> i32 {
- v--;
+int idx3() {
+ v = (v - 1u);
return 3;
}
-fn idx4() -> i32 {
- v--;
+int idx4() {
+ v = (v - 1u);
return 4;
}
-fn idx5() -> i32 {
- v--;
+int idx5() {
+ v = (v - 1u);
return 0;
}
-fn idx6() -> i32 {
- v--;
+int idx6() {
+ v = (v - 1u);
return 2;
}
-fn main() {
- for(buffer[idx1()].a[idx2()][idx3()]--; (v < 10u); buffer[idx4()].a[idx5()][idx6()]--) {
- }
-}
-
-Failed to generate: :73:28 error: binary: %34 is not in scope
- %33:u32 = add %32, %34
- ^^^
-
-:61:7 note: in block
- $B9: { # initializer
- ^^^
-
-:79:9 note: %34 declared here
- %34:u32 = mul %39, 4u
- ^^^^^^^
-
-:112:28 error: binary: %59 is not in scope
- %58:u32 = add %57, %59
- ^^^
-
-:100:7 note: in block
- $B11: { # continuing
- ^^^^
-
-:118:9 note: %59 declared here
- %59:u32 = mul %64, 4u
- ^^^^^^^
-
-note: # Disassembly
-S = struct @align(16) {
- a:array<vec4<i32>, 4> @offset(0)
-}
-
-$B1: { # root
- %buffer:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
- %v:ptr<private, u32, read_write> = var
-}
-
-%idx1 = func():i32 {
- $B2: {
- %4:u32 = load %v
- %5:u32 = sub %4, 1u
- store %v, %5
- ret 1i
- }
-}
-%idx2 = func():i32 {
- $B3: {
- %7:u32 = load %v
- %8:u32 = sub %7, 1u
- store %v, %8
- ret 2i
- }
-}
-%idx3 = func():i32 {
- $B4: {
- %10:u32 = load %v
- %11:u32 = sub %10, 1u
- store %v, %11
- ret 3i
- }
-}
-%idx4 = func():i32 {
- $B5: {
- %13:u32 = load %v
- %14:u32 = sub %13, 1u
- store %v, %14
- ret 4i
- }
-}
-%idx5 = func():i32 {
- $B6: {
- %16:u32 = load %v
- %17:u32 = sub %16, 1u
- store %v, %17
- ret 0i
- }
-}
-%idx6 = func():i32 {
- $B7: {
- %19:u32 = load %v
- %20:u32 = sub %19, 1u
- store %v, %20
- ret 2i
- }
-}
-%main = func():void {
- $B8: {
- loop [i: $B9, b: $B10, c: $B11] { # loop_1
- $B9: { # initializer
- %22:i32 = call %idx1
- %23:i32 = call %idx2
- %24:u32 = convert %22
- %25:u32 = mul %24, 64u
- %26:u32 = convert %23
- %27:u32 = mul %26, 16u
- %28:i32 = call %idx3
- %29:u32 = convert %28
- %30:u32 = mul %29, 4u
- %31:u32 = add 0u, %25
- %32:u32 = add %31, %27
- %33:u32 = add %32, %34
- %35:u32 = add %33, %30
- %36:u32 = %buffer.Load %35
- %37:i32 = bitcast %36
- %38:i32 = sub %37, 1i
- %39:u32 = convert %28
- %34:u32 = mul %39, 4u
- %40:u32 = add 0u, %25
- %41:u32 = add %40, %27
- %42:u32 = add %41, %34
- %43:u32 = bitcast %38
- %44:void = %buffer.Store %42, %43
- next_iteration # -> $B10
+void main() {
+ {
+ int v_1 = idx1();
+ int v_2 = idx2();
+ uint v_3 = (uint(v_1) * 64u);
+ uint v_4 = (uint(v_2) * 16u);
+ int v_5 = idx3();
+ int v_6 = (asint(buffer.Load((((0u + v_3) + v_4) + (uint(v_5) * 4u)))) - 1);
+ uint v_7 = (((0u + v_3) + v_4) + (uint(v_5) * 4u));
+ buffer.Store(v_7, asuint(v_6));
+ while(true) {
+ if ((v < 10u)) {
+ } else {
+ break;
}
- $B10: { # body
- %45:u32 = load %v
- %46:bool = lt %45, 10u
- if %46 [t: $B12, f: $B13] { # if_1
- $B12: { # true
- exit_if # if_1
- }
- $B13: { # false
- exit_loop # loop_1
- }
- }
- continue # -> $B11
+ {
+ int v_8 = idx4();
+ int v_9 = idx5();
+ uint v_10 = (uint(v_8) * 64u);
+ uint v_11 = (uint(v_9) * 16u);
+ int v_12 = idx6();
+ int v_13 = (asint(buffer.Load((((0u + v_10) + v_11) + (uint(v_12) * 4u)))) - 1);
+ uint v_14 = (((0u + v_10) + v_11) + (uint(v_12) * 4u));
+ buffer.Store(v_14, asuint(v_13));
}
- $B11: { # continuing
- %47:i32 = call %idx4
- %48:i32 = call %idx5
- %49:u32 = convert %47
- %50:u32 = mul %49, 64u
- %51:u32 = convert %48
- %52:u32 = mul %51, 16u
- %53:i32 = call %idx6
- %54:u32 = convert %53
- %55:u32 = mul %54, 4u
- %56:u32 = add 0u, %50
- %57:u32 = add %56, %52
- %58:u32 = add %57, %59
- %60:u32 = add %58, %55
- %61:u32 = %buffer.Load %60
- %62:i32 = bitcast %61
- %63:i32 = sub %62, 1i
- %64:u32 = convert %53
- %59:u32 = mul %64, 4u
- %65:u32 = add 0u, %50
- %66:u32 = add %65, %52
- %67:u32 = add %66, %59
- %68:u32 = bitcast %63
- %69:void = %buffer.Store %67, %68
- next_iteration # -> $B10
- }
+ continue;
}
- ret
}
}
-%unused_entry_point = @compute @workgroup_size(1, 1, 1) func():void {
- $B14: {
- ret
- }
+
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
}
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.dxc.hlsl b/test/tint/statements/increment/complex.wgsl.expected.ir.dxc.hlsl
index 12d1d1b..3303817 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.dxc.hlsl
@@ -38,25 +38,23 @@
uint v_3 = (uint(v_1) * 64u);
uint v_4 = (uint(v_2) * 16u);
int v_5 = idx3();
- uint v_6 = (uint(v_5) * 4u);
- int v_7 = (asint(buffer.Load((((0u + v_3) + v_4) + v_6))) + 1);
- uint v_8 = ((((0u + v_3) + v_4) + v_6) + (uint(v_5) * 4u));
- buffer.Store(v_8, asuint(v_7));
+ int v_6 = (asint(buffer.Load((((0u + v_3) + v_4) + (uint(v_5) * 4u)))) + 1);
+ uint v_7 = (((0u + v_3) + v_4) + (uint(v_5) * 4u));
+ buffer.Store(v_7, asuint(v_6));
while(true) {
if ((v < 10u)) {
} else {
break;
}
{
- int v_9 = idx4();
- int v_10 = idx5();
- uint v_11 = (uint(v_9) * 64u);
- uint v_12 = (uint(v_10) * 16u);
- int v_13 = idx6();
- uint v_14 = (uint(v_13) * 4u);
- int v_15 = (asint(buffer.Load((((0u + v_11) + v_12) + v_14))) + 1);
- uint v_16 = ((((0u + v_11) + v_12) + v_14) + (uint(v_13) * 4u));
- buffer.Store(v_16, asuint(v_15));
+ int v_8 = idx4();
+ int v_9 = idx5();
+ uint v_10 = (uint(v_8) * 64u);
+ uint v_11 = (uint(v_9) * 16u);
+ int v_12 = idx6();
+ int v_13 = (asint(buffer.Load((((0u + v_10) + v_11) + (uint(v_12) * 4u)))) + 1);
+ uint v_14 = (((0u + v_10) + v_11) + (uint(v_12) * 4u));
+ buffer.Store(v_14, asuint(v_13));
}
continue;
}
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.fxc.hlsl b/test/tint/statements/increment/complex.wgsl.expected.ir.fxc.hlsl
index 12d1d1b..3303817 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.fxc.hlsl
@@ -38,25 +38,23 @@
uint v_3 = (uint(v_1) * 64u);
uint v_4 = (uint(v_2) * 16u);
int v_5 = idx3();
- uint v_6 = (uint(v_5) * 4u);
- int v_7 = (asint(buffer.Load((((0u + v_3) + v_4) + v_6))) + 1);
- uint v_8 = ((((0u + v_3) + v_4) + v_6) + (uint(v_5) * 4u));
- buffer.Store(v_8, asuint(v_7));
+ int v_6 = (asint(buffer.Load((((0u + v_3) + v_4) + (uint(v_5) * 4u)))) + 1);
+ uint v_7 = (((0u + v_3) + v_4) + (uint(v_5) * 4u));
+ buffer.Store(v_7, asuint(v_6));
while(true) {
if ((v < 10u)) {
} else {
break;
}
{
- int v_9 = idx4();
- int v_10 = idx5();
- uint v_11 = (uint(v_9) * 64u);
- uint v_12 = (uint(v_10) * 16u);
- int v_13 = idx6();
- uint v_14 = (uint(v_13) * 4u);
- int v_15 = (asint(buffer.Load((((0u + v_11) + v_12) + v_14))) + 1);
- uint v_16 = ((((0u + v_11) + v_12) + v_14) + (uint(v_13) * 4u));
- buffer.Store(v_16, asuint(v_15));
+ int v_8 = idx4();
+ int v_9 = idx5();
+ uint v_10 = (uint(v_8) * 64u);
+ uint v_11 = (uint(v_9) * 16u);
+ int v_12 = idx6();
+ int v_13 = (asint(buffer.Load((((0u + v_10) + v_11) + (uint(v_12) * 4u)))) + 1);
+ uint v_14 = (((0u + v_10) + v_11) + (uint(v_12) * 4u));
+ buffer.Store(v_14, asuint(v_13));
}
continue;
}
