VarForDynamicIndex: Skip unevaluated expressions
If an index accessor expression occurs on the RHS of a
constant-evaluated short-circuiting expression, the indices will not
be evaluated if the LHS evaluates to false. Make sure that we skip
these expressions when looking for accesses that need to be hoisted to
vars.
Bug: tint:2038
Change-Id: I34766cc9dd95996cd3047899d7c43b4324b5ad7b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/150260
Auto-Submit: James Price <jrprice@google.com>
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Commit-Queue: David Neto <dneto@google.com>
diff --git a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index.cc b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index.cc
index a560494..e16172f 100644
--- a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index.cc
+++ b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index.cc
@@ -44,8 +44,10 @@
auto* object_expr = access_expr->object;
auto& sem = src->Sem();
- if (sem.GetVal(index_expr)->ConstantValue()) {
- // Index expression resolves to a compile time value.
+ auto stage = sem.GetVal(index_expr)->Stage();
+ if (stage == core::EvaluationStage::kConstant ||
+ stage == core::EvaluationStage::kNotEvaluated) {
+ // Index expression resolves to a compile time value or will not be evaluated.
// As this isn't a dynamic index, we can ignore this.
return true;
}
diff --git a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
index a8f299d..dc27d18 100644
--- a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
+++ b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
@@ -559,5 +559,31 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VarForDynamicIndexTest, ShortCircuitedArrayAccess) {
+ auto* src = R"(
+const foo = (false && (array<f32, 4>()[0] == 0));
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(VarForDynamicIndexTest, ShortCircuitedMatrixAccess) {
+ auto* src = R"(
+const foo = (false && (mat4x4<f32>()[0][0] == 0));
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::spirv::writer
diff --git a/test/tint/bug/tint/2038.wgsl b/test/tint/bug/tint/2038.wgsl
new file mode 100644
index 0000000..bbf9135
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl
@@ -0,0 +1,15 @@
+const foo = (false && (array<f32, 4>()[0] == 0));
+const bar = (false && (mat4x4<f32>()[0][0] == 0));
+
+@group(0) @binding(0)
+var<storage, read_write> output: array<u32, 2>;
+
+@compute @workgroup_size(1)
+fn main() {
+ if (foo) {
+ output[0] = 1;
+ }
+ if (bar) {
+ output[1] = 1;
+ }
+}
\ No newline at end of file
diff --git a/test/tint/bug/tint/2038.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2038.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..3e99de2
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.dxc.hlsl
@@ -0,0 +1,12 @@
+RWByteAddressBuffer output : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ if (false) {
+ output.Store(0u, asuint(1u));
+ }
+ if (false) {
+ output.Store(4u, asuint(1u));
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2038.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2038.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..3e99de2
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.fxc.hlsl
@@ -0,0 +1,12 @@
+RWByteAddressBuffer output : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ if (false) {
+ output.Store(0u, asuint(1u));
+ }
+ if (false) {
+ output.Store(4u, asuint(1u));
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2038.wgsl.expected.glsl b/test/tint/bug/tint/2038.wgsl.expected.glsl
new file mode 100644
index 0000000..ad36cb6
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.glsl
@@ -0,0 +1,20 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ uint inner[2];
+} tint_symbol;
+
+void tint_symbol_1() {
+ if (false) {
+ tint_symbol.inner[0] = 1u;
+ }
+ if (false) {
+ tint_symbol.inner[1] = 1u;
+ }
+}
+
+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/2038.wgsl.expected.msl b/test/tint/bug/tint/2038.wgsl.expected.msl
new file mode 100644
index 0000000..a091f37
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#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];
+};
+
+kernel void tint_symbol(device tint_array<uint, 2>* tint_symbol_1 [[buffer(0)]]) {
+ if (false) {
+ (*(tint_symbol_1))[0] = 1u;
+ }
+ if (false) {
+ (*(tint_symbol_1))[1] = 1u;
+ }
+ return;
+}
+
diff --git a/test/tint/bug/tint/2038.wgsl.expected.spvasm b/test/tint/bug/tint/2038.wgsl.expected.spvasm
new file mode 100644
index 0000000..fd316e3
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.spvasm
@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %output_block "output_block"
+ OpMemberName %output_block 0 "inner"
+ OpName %output "output"
+ OpName %main "main"
+ OpDecorate %output_block Block
+ OpMemberDecorate %output_block 0 Offset 0
+ OpDecorate %_arr_uint_uint_2 ArrayStride 4
+ OpDecorate %output DescriptorSet 0
+ OpDecorate %output Binding 0
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
+%output_block = OpTypeStruct %_arr_uint_uint_2
+%_ptr_StorageBuffer_output_block = OpTypePointer StorageBuffer %output_block
+ %output = OpVariable %_ptr_StorageBuffer_output_block StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %bool = OpTypeBool
+ %12 = OpConstantNull %bool
+ %uint_0 = OpConstant %uint 0
+ %int = OpTypeInt 32 1
+ %17 = OpConstantNull %int
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %uint_1 = OpConstant %uint 1
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ OpSelectionMerge %13 None
+ OpBranchConditional %12 %14 %13
+ %14 = OpLabel
+ %19 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %17
+ OpStore %19 %uint_1
+ OpBranch %13
+ %13 = OpLabel
+ OpSelectionMerge %21 None
+ OpBranchConditional %12 %22 %21
+ %22 = OpLabel
+ %24 = OpAccessChain %_ptr_StorageBuffer_uint %output %uint_0 %int_1
+ OpStore %24 %uint_1
+ OpBranch %21
+ %21 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2038.wgsl.expected.wgsl b/test/tint/bug/tint/2038.wgsl.expected.wgsl
new file mode 100644
index 0000000..001d68b
--- /dev/null
+++ b/test/tint/bug/tint/2038.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+const foo = (false && (array<f32, 4>()[0] == 0));
+
+const bar = (false && (mat4x4<f32>()[0][0] == 0));
+
+@group(0) @binding(0) var<storage, read_write> output : array<u32, 2>;
+
+@compute @workgroup_size(1)
+fn main() {
+ if (foo) {
+ output[0] = 1;
+ }
+ if (bar) {
+ output[1] = 1;
+ }
+}
