From 5716611bf0ad32be0a7addad8fac630c3cdbd4f6 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 21 Jul 2022 15:25:35 +0000 Subject: [PATCH] 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 Kokoro: Kokoro Commit-Queue: Ben Clayton --- src/tint/resolver/materialize_test.cc | 56 ++++++++++++++++++- src/tint/resolver/resolver.cc | 8 ++- test/tint/bug/chromium/1345468.wgsl | 8 +++ .../bug/chromium/1345468.wgsl.expected.glsl | 12 ++++ .../bug/chromium/1345468.wgsl.expected.hlsl | 10 ++++ .../bug/chromium/1345468.wgsl.expected.msl | 9 +++ .../bug/chromium/1345468.wgsl.expected.spvasm | 55 ++++++++++++++++++ .../bug/chromium/1345468.wgsl.expected.wgsl | 7 +++ 8 files changed, 161 insertions(+), 4 deletions(-) create mode 100644 test/tint/bug/chromium/1345468.wgsl create mode 100644 test/tint/bug/chromium/1345468.wgsl.expected.glsl create mode 100644 test/tint/bug/chromium/1345468.wgsl.expected.hlsl create mode 100644 test/tint/bug/chromium/1345468.wgsl.expected.msl create mode 100644 test/tint/bug/chromium/1345468.wgsl.expected.spvasm create mode 100644 test/tint/bug/chromium/1345468.wgsl.expected.wgsl diff --git a/src/tint/resolver/materialize_test.cc b/src/tint/resolver/materialize_test.cc index 1b935eef0e..56492b099f 100644 --- a/src/tint/resolver/materialize_test.cc +++ b/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 << ""; } @@ -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(AFloat(-kSubnormalF16), -kSubnormalF16), */ // }))); +INSTANTIATE_TEST_SUITE_P( + MaterializeVectorRuntimeIndex, + MaterializeAbstractNumericToConcreteType, + testing::Combine(testing::Values(Expectation::kMaterialize), + testing::Values(Method::kRuntimeIndex), + testing::ValuesIn(std::vector{ + Types(0_a, 0.0), // + Types(1_a, 1.0), // + Types(-1_a, -1.0), // + Types(AInt(kHighestI32), kHighestI32), // + Types(AInt(kLowestI32), kLowestI32), // + Types(0.0_a, 0.0), // + Types(1.0_a, 1.0), // + Types(-1.0_a, -1.0), // + Types(AFloat(kHighestF32), kHighestF32), // + Types(AFloat(kLowestF32), kLowestF32), // + Types(AFloat(kPiF32), kPiF64), // + Types(AFloat(kSubnormalF32), kSubnormalF32), // + Types(AFloat(-kSubnormalF32), -kSubnormalF32), // + }))); + INSTANTIATE_TEST_SUITE_P( MaterializeMatrix, MaterializeAbstractNumericToConcreteType, @@ -535,6 +565,22 @@ INSTANTIATE_TEST_SUITE_P( /* Types(AFloat(-kSubnormalF16), -kSubnormalF16), */ // }))); +INSTANTIATE_TEST_SUITE_P( + MaterializeMatrixRuntimeIndex, + MaterializeAbstractNumericToConcreteType, + testing::Combine(testing::Values(Expectation::kMaterialize), + testing::Values(Method::kRuntimeIndex), + testing::ValuesIn(std::vector{ + Types(0.0_a, 0.0), // + Types(1.0_a, 1.0), // + Types(-1.0_a, -1.0), // + Types(AFloat(kHighestF32), kHighestF32), // + Types(AFloat(kLowestF32), kLowestF32), // + Types(AFloat(kPiF32), kPiF64), // + Types(AFloat(kSubnormalF32), kSubnormalF32), // + Types(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 diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc index 611f240f3c..db9851cf52 100644 --- a/src/tint/resolver/resolver.cc +++ b/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( diff --git a/test/tint/bug/chromium/1345468.wgsl b/test/tint/bug/chromium/1345468.wgsl new file mode 100644 index 0000000000..97369e9b89 --- /dev/null +++ b/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 +} diff --git a/test/tint/bug/chromium/1345468.wgsl.expected.glsl b/test/tint/bug/chromium/1345468.wgsl.expected.glsl new file mode 100644 index 0000000000..e915c629c0 --- /dev/null +++ b/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]; +} + diff --git a/test/tint/bug/chromium/1345468.wgsl.expected.hlsl b/test/tint/bug/chromium/1345468.wgsl.expected.hlsl new file mode 100644 index 0000000000..6b3556b9f5 --- /dev/null +++ b/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]; +} diff --git a/test/tint/bug/chromium/1345468.wgsl.expected.msl b/test/tint/bug/chromium/1345468.wgsl.expected.msl new file mode 100644 index 0000000000..c0c1a74021 --- /dev/null +++ b/test/tint/bug/chromium/1345468.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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]; +} + diff --git a/test/tint/bug/chromium/1345468.wgsl.expected.spvasm b/test/tint/bug/chromium/1345468.wgsl.expected.spvasm new file mode 100644 index 0000000000..17453d9e34 --- /dev/null +++ b/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 diff --git a/test/tint/bug/chromium/1345468.wgsl.expected.wgsl b/test/tint/bug/chromium/1345468.wgsl.expected.wgsl new file mode 100644 index 0000000000..014a19b495 --- /dev/null +++ b/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]; +}