From 3fe1bd3715f494d7483a6cc69a151236e224a1a9 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Fri, 28 Jan 2022 16:49:46 +0000 Subject: [PATCH] transform: Fix CalculateArrayLength for arrays The transform was not correctly inserting the intrinsic call after array element types. Fixed: chromium:1290107 Change-Id: I7199d1846cb98305d789cf0bc362eb5872d9b917 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78542 Kokoro: Kokoro Reviewed-by: James Price --- src/transform/calculate_array_length.cc | 9 ++- src/transform/calculate_array_length_test.cc | 76 ++++++++++++++++++- test/bug/chromium/1290107.wgsl | 10 +++ test/bug/chromium/1290107.wgsl.expected.glsl | 32 ++++++++ test/bug/chromium/1290107.wgsl.expected.hlsl | 10 +++ test/bug/chromium/1290107.wgsl.expected.msl | 15 ++++ .../bug/chromium/1290107.wgsl.expected.spvasm | 36 +++++++++ test/bug/chromium/1290107.wgsl.expected.wgsl | 10 +++ 8 files changed, 196 insertions(+), 2 deletions(-) create mode 100644 test/bug/chromium/1290107.wgsl create mode 100644 test/bug/chromium/1290107.wgsl.expected.glsl create mode 100644 test/bug/chromium/1290107.wgsl.expected.hlsl create mode 100644 test/bug/chromium/1290107.wgsl.expected.msl create mode 100644 test/bug/chromium/1290107.wgsl.expected.spvasm create mode 100644 test/bug/chromium/1290107.wgsl.expected.wgsl diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc index c54c3e8aaf..b365f06382 100644 --- a/src/transform/calculate_array_length.cc +++ b/src/transform/calculate_array_length.cc @@ -119,7 +119,14 @@ void CalculateArrayLength::Run(CloneContext& ctx, ctx.dst->ASTNodes().Create(ctx.dst->ID()), }, ast::DecorationList{}); - if (auto* str = buffer_type->As()) { + // Insert the intrinsic function after the structure or array structure + // element type. TODO(crbug.com/tint/1266): Once we allow out-of-order + // declarations, this can be simplified. + auto* insert_after = buffer_type; + while (auto* arr = insert_after->As()) { + insert_after = arr->ElemType(); + } + if (auto* str = insert_after->As()) { ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(), func); } else { diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc index cebb27647e..fcf70010d3 100644 --- a/src/transform/calculate_array_length_test.cc +++ b/src/transform/calculate_array_length_test.cc @@ -65,7 +65,7 @@ fn main() { EXPECT_TRUE(ShouldRun(src)); } -TEST_F(CalculateArrayLengthTest, Basic) { +TEST_F(CalculateArrayLengthTest, BasicArray) { auto* src = R"( @group(0) @binding(0) var sb : array; @@ -135,6 +135,80 @@ fn main() { EXPECT_EQ(expect, str(got)); } +TEST_F(CalculateArrayLengthTest, ArrayOfStruct) { + auto* src = R"( +struct S { + f : f32; +} + +@group(0) @binding(0) var arr : array; + +@stage(compute) @workgroup_size(1) +fn main() { + let len = arrayLength(&arr); +} +)"; + auto* expect = R"( +struct S { + f : f32; +} + +@internal(intrinsic_buffer_size) +fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array, result : ptr) + +@group(0) @binding(0) var arr : array; + +@stage(compute) @workgroup_size(1) +fn main() { + var tint_symbol_1 : u32 = 0u; + tint_symbol(arr, &(tint_symbol_1)); + let tint_symbol_2 : u32 = (tint_symbol_1 / 4u); + let len = tint_symbol_2; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(CalculateArrayLengthTest, ArrayOfArrayOfStruct) { + auto* src = R"( +struct S { + f : f32; +} + +@group(0) @binding(0) var arr : array>; + +@stage(compute) @workgroup_size(1) +fn main() { + let len = arrayLength(&arr); +} +)"; + auto* expect = R"( +struct S { + f : f32; +} + +@internal(intrinsic_buffer_size) +fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array>, result : ptr) + +@group(0) @binding(0) var arr : array>; + +@stage(compute) @workgroup_size(1) +fn main() { + var tint_symbol_1 : u32 = 0u; + tint_symbol(arr, &(tint_symbol_1)); + let tint_symbol_2 : u32 = (tint_symbol_1 / 16u); + let len = tint_symbol_2; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(CalculateArrayLengthTest, InSameBlock) { auto* src = R"( @group(0) @binding(0) var sb : array;; diff --git a/test/bug/chromium/1290107.wgsl b/test/bug/chromium/1290107.wgsl new file mode 100644 index 0000000000..c2e477865e --- /dev/null +++ b/test/bug/chromium/1290107.wgsl @@ -0,0 +1,10 @@ +struct S { + f : f32; +} + +@group(0) @binding(0) var arr : array; + +@stage(compute) @workgroup_size(1) +fn main() { + let len = arrayLength(&arr); +} diff --git a/test/bug/chromium/1290107.wgsl.expected.glsl b/test/bug/chromium/1290107.wgsl.expected.glsl new file mode 100644 index 0000000000..debb813788 --- /dev/null +++ b/test/bug/chromium/1290107.wgsl.expected.glsl @@ -0,0 +1,32 @@ +SKIP: FAILED + +#version 310 es +precision mediump float; + +struct S { + float f; +}; + +layout(binding = 0) buffer arr_block_1 { + S inner[]; +} arr; +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void tint_symbol() { + uint tint_symbol_2 = 0u; + arr.inner.GetDimensions(tint_symbol_2); + uint tint_symbol_3 = (tint_symbol_2 / 4u); + uint len = tint_symbol_3; + return; +} + +void main() { + tint_symbol(); +} + +Error parsing GLSL shader: +ERROR: 0:14: '.' : cannot apply to an array: GetDimensions +ERROR: 0:14: '' : compilation terminated +ERROR: 2 compilation errors. No code generated. + + + diff --git a/test/bug/chromium/1290107.wgsl.expected.hlsl b/test/bug/chromium/1290107.wgsl.expected.hlsl new file mode 100644 index 0000000000..95e8cb7e13 --- /dev/null +++ b/test/bug/chromium/1290107.wgsl.expected.hlsl @@ -0,0 +1,10 @@ +ByteAddressBuffer arr : register(t0, space0); + +[numthreads(1, 1, 1)] +void main() { + uint tint_symbol_1 = 0u; + arr.GetDimensions(tint_symbol_1); + const uint tint_symbol_2 = (tint_symbol_1 / 4u); + const uint len = tint_symbol_2; + return; +} diff --git a/test/bug/chromium/1290107.wgsl.expected.msl b/test/bug/chromium/1290107.wgsl.expected.msl new file mode 100644 index 0000000000..425519536e --- /dev/null +++ b/test/bug/chromium/1290107.wgsl.expected.msl @@ -0,0 +1,15 @@ +#include + +using namespace metal; +struct tint_symbol_1 { + /* 0x0000 */ uint4 buffer_size[1]; +}; +struct S { + float f; +}; + +kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) { + uint const len = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u); + return; +} + diff --git a/test/bug/chromium/1290107.wgsl.expected.spvasm b/test/bug/chromium/1290107.wgsl.expected.spvasm new file mode 100644 index 0000000000..1c835ab465 --- /dev/null +++ b/test/bug/chromium/1290107.wgsl.expected.spvasm @@ -0,0 +1,36 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 13 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %arr_block "arr_block" + OpMemberName %arr_block 0 "inner" + OpName %S "S" + OpMemberName %S 0 "f" + OpName %arr "arr" + OpName %main "main" + OpDecorate %arr_block Block + OpMemberDecorate %arr_block 0 Offset 0 + OpMemberDecorate %S 0 Offset 0 + OpDecorate %_runtimearr_S ArrayStride 4 + OpDecorate %arr NonWritable + OpDecorate %arr DescriptorSet 0 + OpDecorate %arr Binding 0 + %float = OpTypeFloat 32 + %S = OpTypeStruct %float +%_runtimearr_S = OpTypeRuntimeArray %S + %arr_block = OpTypeStruct %_runtimearr_S +%_ptr_StorageBuffer_arr_block = OpTypePointer StorageBuffer %arr_block + %arr = OpVariable %_ptr_StorageBuffer_arr_block StorageBuffer + %void = OpTypeVoid + %7 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %main = OpFunction %void None %7 + %10 = OpLabel + %11 = OpArrayLength %uint %arr 0 + OpReturn + OpFunctionEnd diff --git a/test/bug/chromium/1290107.wgsl.expected.wgsl b/test/bug/chromium/1290107.wgsl.expected.wgsl new file mode 100644 index 0000000000..4041cefffa --- /dev/null +++ b/test/bug/chromium/1290107.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +struct S { + f : f32; +} + +@group(0) @binding(0) var arr : array; + +@stage(compute) @workgroup_size(1) +fn main() { + let len = arrayLength(&(arr)); +}