tint/resolver: Materialize objects when indexed with non-const index

If an abstract-vector or abstract-matrix is indexed with a non-constant index expression, then the resulting value is non-constant, and so cannot be abstract.

In this situation the materialization cannot be done post-index, so materialization must happen on the object before indexing.

Bug: chromium:1345468
Change-Id: I9f29dc40301779a7ff8f173724374bd845a3a5b9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/96684
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>

src/tint/resolver/materialize_test.cc

@@ -200,7 +200,10 @@ enum class Method {
     // @workgroup_size(target_expr, abstract_expr, 123)
     // @compute
     // fn f() {}
-    kWorkgroupSize
+    kWorkgroupSize,
+
+    // abstract_expr[runtime-index]
+    kRuntimeIndex,
 };
 
 static std::ostream& operator<<(std::ostream& o, Method m) {
@@ -235,6 +238,8 @@ static std::ostream& operator<<(std::ostream& o, Method m) {
             return o << "switch-case-with-abstract";
         case Method::kWorkgroupSize:
             return o << "workgroup-size";
+        case Method::kRuntimeIndex:
+            return o << "dynamic-index";
     }
     return o << "<unknown>";
 }
@@ -370,6 +375,10 @@ TEST_P(MaterializeAbstractNumericToConcreteType, Test) {
                  {WorkgroupSize(target_expr(), abstract_expr, Expr(123_a)),
                   Stage(ast::PipelineStage::kCompute)});
             break;
+        case Method::kRuntimeIndex:
+            auto* runtime_index = Var("runtime_index", nullptr, Expr(1_i));
+            WrapInFunction(runtime_index, IndexAccessor(abstract_expr, runtime_index));
+            break;
     }
 
     switch (expectation) {
@@ -511,6 +520,27 @@ INSTANTIATE_TEST_SUITE_P(
                          /* Types<f16V, AFloatV>(AFloat(-kSubnormalF16), -kSubnormalF16), */  //
                      })));
 
+INSTANTIATE_TEST_SUITE_P(
+    MaterializeVectorRuntimeIndex,
+    MaterializeAbstractNumericToConcreteType,
+    testing::Combine(testing::Values(Expectation::kMaterialize),
+                     testing::Values(Method::kRuntimeIndex),
+                     testing::ValuesIn(std::vector<Data>{
+                         Types<i32V, AIntV>(0_a, 0.0),                                  //
+                         Types<i32V, AIntV>(1_a, 1.0),                                  //
+                         Types<i32V, AIntV>(-1_a, -1.0),                                //
+                         Types<i32V, AIntV>(AInt(kHighestI32), kHighestI32),            //
+                         Types<i32V, AIntV>(AInt(kLowestI32), kLowestI32),              //
+                         Types<f32V, AFloatV>(0.0_a, 0.0),                              //
+                         Types<f32V, AFloatV>(1.0_a, 1.0),                              //
+                         Types<f32V, AFloatV>(-1.0_a, -1.0),                            //
+                         Types<f32V, AFloatV>(AFloat(kHighestF32), kHighestF32),        //
+                         Types<f32V, AFloatV>(AFloat(kLowestF32), kLowestF32),          //
+                         Types<f32V, AFloatV>(AFloat(kPiF32), kPiF64),                  //
+                         Types<f32V, AFloatV>(AFloat(kSubnormalF32), kSubnormalF32),    //
+                         Types<f32V, AFloatV>(AFloat(-kSubnormalF32), -kSubnormalF32),  //
+                     })));
+
 INSTANTIATE_TEST_SUITE_P(
     MaterializeMatrix,
     MaterializeAbstractNumericToConcreteType,
@@ -535,6 +565,22 @@ INSTANTIATE_TEST_SUITE_P(
                          /* Types<f16M, AFloatM>(AFloat(-kSubnormalF16), -kSubnormalF16), */  //
                      })));
 
+INSTANTIATE_TEST_SUITE_P(
+    MaterializeMatrixRuntimeIndex,
+    MaterializeAbstractNumericToConcreteType,
+    testing::Combine(testing::Values(Expectation::kMaterialize),
+                     testing::Values(Method::kRuntimeIndex),
+                     testing::ValuesIn(std::vector<Data>{
+                         Types<f32M, AFloatM>(0.0_a, 0.0),                              //
+                         Types<f32M, AFloatM>(1.0_a, 1.0),                              //
+                         Types<f32M, AFloatM>(-1.0_a, -1.0),                            //
+                         Types<f32M, AFloatM>(AFloat(kHighestF32), kHighestF32),        //
+                         Types<f32M, AFloatM>(AFloat(kLowestF32), kLowestF32),          //
+                         Types<f32M, AFloatM>(AFloat(kPiF32), kPiF64),                  //
+                         Types<f32M, AFloatM>(AFloat(kSubnormalF32), kSubnormalF32),    //
+                         Types<f32M, AFloatM>(AFloat(-kSubnormalF32), -kSubnormalF32),  //
+                     })));
+
 INSTANTIATE_TEST_SUITE_P(MaterializeSwitch,
                          MaterializeAbstractNumericToConcreteType,
                          testing::Combine(testing::Values(Expectation::kMaterialize),
@@ -998,12 +1044,14 @@ INSTANTIATE_TEST_SUITE_P(ArrayLengthValueCannotBeRepresented,
 
 }  // namespace materialize_abstract_numeric_to_default_type
 
+namespace materialize_abstract_numeric_to_unrelated_type {
+
 using MaterializeAbstractNumericToUnrelatedType = resolver::ResolverTest;
 
 TEST_F(MaterializeAbstractNumericToUnrelatedType, AIntToStructVarCtor) {
     Structure("S", {Member("a", ty.i32())});
     WrapInFunction(Decl(Var("v", ty.type_name("S"), Expr(Source{{12, 34}}, 1_a))));
-    ASSERT_FALSE(r()->Resolve());
+    EXPECT_FALSE(r()->Resolve());
     EXPECT_THAT(
         r()->error(),
         testing::HasSubstr("error: cannot convert value of type 'abstract-int' to type 'S'"));
@@ -1012,11 +1060,13 @@ TEST_F(MaterializeAbstractNumericToUnrelatedType, AIntToStructVarCtor) {
 TEST_F(MaterializeAbstractNumericToUnrelatedType, AIntToStructLetCtor) {
     Structure("S", {Member("a", ty.i32())});
     WrapInFunction(Decl(Let("v", ty.type_name("S"), Expr(Source{{12, 34}}, 1_a))));
-    ASSERT_FALSE(r()->Resolve());
+    EXPECT_FALSE(r()->Resolve());
     EXPECT_THAT(
         r()->error(),
         testing::HasSubstr("error: cannot convert value of type 'abstract-int' to type 'S'"));
 }
 
+}  // namespace materialize_abstract_numeric_to_unrelated_type
+
 }  // namespace
 }  // namespace tint::resolver

src/tint/resolver/resolver.cc

@@ -1399,7 +1399,13 @@ sem::Expression* Resolver::IndexAccessor(const ast::IndexAccessorExpression* exp
     if (!idx) {
         return nullptr;
     }
-    auto* obj = sem_.Get(expr->object);
+    const auto* obj = sem_.Get(expr->object);
+    if (idx->Stage() != sem::EvaluationStage::kConstant) {
+        // If the index is non-constant, then the resulting expression is non-constant, so we'll
+        // have to materialize the object. For example, consider:
+        //     vec2(1, 2)[runtime-index]
+        obj = Materialize(obj);
+    }
     auto* obj_raw_ty = obj->Type();
     auto* obj_ty = obj_raw_ty->UnwrapRef();
     auto* ty = Switch(

test/tint/bug/chromium/1345468.wgsl

@@ -0,0 +1,8 @@
+
+fn f(){
+    const m = mat4x2(0, 0, 0, 0, 4., 0, 0, 0); // abstract matrix
+    const v = vec2(0, 1);                      // abstract vector
+    var i = 1;                                 // runtime-evaluated index
+    var a = m[i];                              // materialize m before index
+    var b = v[i];                              // materialize v before index
+}

test/tint/bug/chromium/1345468.wgsl.expected.glsl

@@ -0,0 +1,12 @@
+#version 310 es
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void unused_entry_point() {
+  return;
+}
+void f() {
+  int i = 1;
+  vec2 a = mat4x2(vec2(0.0f), vec2(0.0f), vec2(4.0f, 0.0f), vec2(0.0f))[i];
+  int b = ivec2(0, 1)[i];
+}
+

test/tint/bug/chromium/1345468.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  int i = 1;
+  float2 a = float4x2((0.0f).xx, (0.0f).xx, float2(4.0f, 0.0f), (0.0f).xx)[i];
+  int b = int2(0, 1)[i];
+}

test/tint/bug/chromium/1345468.wgsl.expected.msl

@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void f() {
+  int i = 1;
+  float2 a = float4x2(float2(0.0f), float2(0.0f), float2(4.0f, 0.0f), float2(0.0f))[i];
+  int b = int2(0, 1)[i];
+}
+

test/tint/bug/chromium/1345468.wgsl.expected.spvasm

@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpName %var_for_index "var_for_index"
+               OpName %a "a"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %11 = OpConstantNull %int
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat4v2float = OpTypeMatrix %v2float 4
+         %15 = OpConstantNull %v2float
+    %float_4 = OpConstant %float 4
+         %17 = OpConstantNull %float
+         %18 = OpConstantComposite %v2float %float_4 %17
+         %19 = OpConstantComposite %mat4v2float %15 %15 %18 %15
+%_ptr_Function_mat4v2float = OpTypePointer Function %mat4v2float
+         %22 = OpConstantNull %mat4v2float
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+      %v2int = OpTypeVector %int 2
+         %29 = OpConstantComposite %v2int %11 %int_1
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %11
+%var_for_index = OpVariable %_ptr_Function_mat4v2float Function %22
+          %a = OpVariable %_ptr_Function_v2float Function %15
+          %b = OpVariable %_ptr_Function_int Function %11
+               OpStore %i %int_1
+               OpStore %var_for_index %19
+         %23 = OpLoad %int %i
+         %25 = OpAccessChain %_ptr_Function_v2float %var_for_index %23
+         %26 = OpLoad %v2float %25
+               OpStore %a %26
+         %30 = OpLoad %int %i
+         %31 = OpVectorExtractDynamic %int %29 %30
+               OpStore %b %31
+               OpReturn
+               OpFunctionEnd

test/tint/bug/chromium/1345468.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+fn f() {
+  const m = mat4x2(0, 0, 0, 0, 4.0, 0, 0, 0);
+  const v = vec2(0, 1);
+  var i = 1;
+  var a = m[i];
+  var b = v[i];
+}