From de2b7db2447018e226838a8152d86a235c0c9f16 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Wed, 14 Jul 2021 17:28:01 +0000 Subject: [PATCH] Make ArrayLengthFromUniform transform emit a valid UBO The UBO must have a stride that is a multiple of 16 bytes. Note that this change was part of https://dawn-review.googlesource.com/c/tint/+/56780 but the CL was reverted because it broke Dawn. This CL relands part of the change, and adds the macro TINT_EXPECTS_UBOS_TO_BE_MULTIPLE_OF_16 so that Dawn can conditionally compile against it. Bug: tint:984 Bug: tint:643 Change-Id: I303b3fe81ff97c4933c489736d5d5432a59ce9b7 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57921 Kokoro: Kokoro Commit-Queue: Antonio Maiorano Reviewed-by: Ben Clayton --- include/tint/tint.h | 3 + src/transform/array_length_from_uniform.cc | 28 +++++--- .../array_length_from_uniform_test.cc | 71 ++++++++++++++++--- .../complex_via_let.wgsl.expected.msl | 4 +- .../arrayLength/deprecated.wgsl.expected.msl | 6 +- .../arrayLength/simple.wgsl.expected.msl | 4 +- .../arrayLength/via_let.wgsl.expected.msl | 4 +- .../via_let_complex.wgsl.expected.msl | 4 +- .../gen/arrayLength/1588cd.wgsl.expected.msl | 4 +- .../gen/arrayLength/61b1c7.wgsl.expected.msl | 4 +- .../gen/arrayLength/a0f5ca.wgsl.expected.msl | 4 +- .../gen/arrayLength/cdd123.wgsl.expected.msl | 4 +- .../gen/arrayLength/cfca0a.wgsl.expected.msl | 4 +- .../gen/arrayLength/eb510f.wgsl.expected.msl | 4 +- 14 files changed, 106 insertions(+), 42 deletions(-) diff --git a/include/tint/tint.h b/include/tint/tint.h index ac1c8b3563..9d28687a05 100644 --- a/include/tint/tint.h +++ b/include/tint/tint.h @@ -59,4 +59,7 @@ #include "src/writer/hlsl/generator.h" #endif // TINT_BUILD_HLSL_WRITER +// TODO(crbug/984): Remove once Dawn builds with this flag +#define TINT_EXPECTS_UBOS_TO_BE_MULTIPLE_OF_16 + #endif // INCLUDE_TINT_TINT_H_ diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc index bb182c9380..35fbcb999f 100644 --- a/src/transform/array_length_from_uniform.cc +++ b/src/transform/array_length_from_uniform.cc @@ -67,11 +67,16 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx, ast::Variable* buffer_size_ubo = nullptr; auto get_ubo = [&]() { if (!buffer_size_ubo) { + // Emit an array, N>, where N is 1/4 number of elements. + // We do this because UBOs require an element stride that is 16-byte + // aligned. auto* buffer_size_struct = ctx.dst->Structure( ctx.dst->Sym(), {ctx.dst->Member( kBufferSizeMemberName, - ctx.dst->ty.array(ctx.dst->ty.u32(), max_buffer_size_index + 1))}, + ctx.dst->ty.array(ctx.dst->ty.vec4(ctx.dst->ty.u32()), + (max_buffer_size_index / 4) + 1))}, + ast::DecorationList{ctx.dst->create()}); buffer_size_ubo = ctx.dst->Global( ctx.dst->Sym(), ctx.dst->ty.Of(buffer_size_struct), @@ -99,18 +104,20 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx, // Get the storage buffer that contains the runtime array. // We assume that the argument to `arrayLength` has the form - // `&resource.array`, which requires that `InlinePointerLets` and `Simplify` - // have been run before this transform. + // `&resource.array`, which requires that `InlinePointerLets` and + // `Simplify` have been run before this transform. auto* param = call_expr->params()[0]->As(); if (!param || param->op() != ast::UnaryOp::kAddressOf) { TINT_ICE(Transform, ctx.dst->Diagnostics()) - << "expected form of arrayLength argument to be &resource.array"; + << "expected form of arrayLength argument to be " + "&resource.array"; break; } auto* accessor = param->expr()->As(); if (!accessor) { TINT_ICE(Transform, ctx.dst->Diagnostics()) - << "expected form of arrayLength argument to be &resource.array"; + << "expected form of arrayLength argument to be " + "&resource.array"; break; } auto* storage_buffer_expr = accessor->structure(); @@ -118,7 +125,8 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx, sem.Get(storage_buffer_expr)->As(); if (!storage_buffer_sem) { TINT_ICE(Transform, ctx.dst->Diagnostics()) - << "expected form of arrayLength argument to be &resource.array"; + << "expected form of arrayLength argument to be " + "&resource.array"; break; } @@ -135,9 +143,13 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx, } // Load the total storage buffer size from the UBO. - auto* total_storage_buffer_size = ctx.dst->IndexAccessor( + uint32_t array_index = idx_itr->second / 4; + auto* vec_expr = ctx.dst->IndexAccessor( ctx.dst->MemberAccessor(get_ubo()->symbol(), kBufferSizeMemberName), - idx_itr->second); + array_index); + uint32_t vec_index = idx_itr->second % 4; + auto* total_storage_buffer_size = + ctx.dst->IndexAccessor(vec_expr, vec_index); // Calculate actual array length // total_storage_buffer_size - array_offset diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc index a173de8ab4..6ab39ee055 100644 --- a/src/transform/array_length_from_uniform_test.cc +++ b/src/transform/array_length_from_uniform_test.cc @@ -81,7 +81,7 @@ fn main() { auto* expect = R"( [[block]] struct tint_symbol { - buffer_size : array; + buffer_size : array, 1>; }; [[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; @@ -96,7 +96,7 @@ struct SB { [[stage(compute), workgroup_size(1)]] fn main() { - var len : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u); + var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u); } )"; @@ -134,7 +134,7 @@ fn main() { auto* expect = R"( [[block]] struct tint_symbol { - buffer_size : array; + buffer_size : array, 1>; }; [[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; @@ -150,7 +150,7 @@ struct SB { [[stage(compute), workgroup_size(1)]] fn main() { - var len : u32 = ((tint_symbol_1.buffer_size[0u] - 8u) / 64u); + var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u); } )"; @@ -175,29 +175,48 @@ struct SB1 { x : i32; arr1 : array; }; - [[block]] struct SB2 { x : i32; arr2 : array>; }; +[[block]] +struct SB3 { + x : i32; + arr3 : array>; +}; +[[block]] +struct SB4 { + x : i32; + arr4 : array>; +}; +[[block]] +struct SB5 { + x : i32; + arr5 : array>; +}; [[group(0), binding(2)]] var sb1 : SB1; - [[group(1), binding(2)]] var sb2 : SB2; +[[group(2), binding(2)]] var sb3 : SB3; +[[group(3), binding(2)]] var sb4 : SB4; +[[group(4), binding(2)]] var sb5 : SB5; [[stage(compute), workgroup_size(1)]] fn main() { var len1 : u32 = arrayLength(&(sb1.arr1)); var len2 : u32 = arrayLength(&(sb2.arr2)); - var x : u32 = (len1 + len2); + var len3 : u32 = arrayLength(&(sb3.arr3)); + var len4 : u32 = arrayLength(&(sb4.arr4)); + var len5 : u32 = arrayLength(&(sb5.arr5)); + var x : u32 = (len1 + len2 + len3 + len4 + len5); } )"; auto* expect = R"( [[block]] struct tint_symbol { - buffer_size : array; + buffer_size : array, 2>; }; [[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; @@ -214,21 +233,51 @@ struct SB2 { arr2 : array>; }; +[[block]] +struct SB3 { + x : i32; + arr3 : array>; +}; + +[[block]] +struct SB4 { + x : i32; + arr4 : array>; +}; + +[[block]] +struct SB5 { + x : i32; + arr5 : array>; +}; + [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; +[[group(2), binding(2)]] var sb3 : SB3; + +[[group(3), binding(2)]] var sb4 : SB4; + +[[group(4), binding(2)]] var sb5 : SB5; + [[stage(compute), workgroup_size(1)]] fn main() { - var len1 : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u); - var len2 : u32 = ((tint_symbol_1.buffer_size[1u] - 16u) / 16u); - var x : u32 = (len1 + len2); + var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u); + var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u); + var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u); + var len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u); + var len5 : u32 = ((tint_symbol_1.buffer_size[1u][0u] - 16u) / 16u); + var x : u32 = ((((len1 + len2) + len3) + len4) + len5); } )"; ArrayLengthFromUniform::Config cfg({0, 30u}); cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2u}, 0); cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{1u, 2u}, 1); + cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{2u, 2u}, 2); + cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{3u, 2u}, 3); + cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{4u, 2u}, 4); DataMap data; data.Add(std::move(cfg)); diff --git a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl +++ b/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl @@ -2,14 +2,14 @@ using namespace metal; struct tint_symbol_1 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct S { /* 0x0000 */ int a[1]; }; kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { - uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); + uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); return; } diff --git a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl index c86fc59bd4..40f5a450e5 100644 --- a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl +++ b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl @@ -2,15 +2,15 @@ using namespace metal; struct tint_symbol_1 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct S { /* 0x0000 */ int a[1]; }; kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { - uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); - uint const l2 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); + uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); + uint const l2 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); return; } diff --git a/test/intrinsics/arrayLength/simple.wgsl.expected.msl b/test/intrinsics/arrayLength/simple.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/simple.wgsl.expected.msl +++ b/test/intrinsics/arrayLength/simple.wgsl.expected.msl @@ -2,14 +2,14 @@ using namespace metal; struct tint_symbol_1 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct S { /* 0x0000 */ int a[1]; }; kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { - uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); + uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); return; } diff --git a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl +++ b/test/intrinsics/arrayLength/via_let.wgsl.expected.msl @@ -2,14 +2,14 @@ using namespace metal; struct tint_symbol_1 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct S { /* 0x0000 */ int a[1]; }; kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { - uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); + uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); return; } diff --git a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl +++ b/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl @@ -2,14 +2,14 @@ using namespace metal; struct tint_symbol_1 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct S { /* 0x0000 */ int a[1]; }; kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { - uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u); + uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u); return; } diff --git a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl index 4f53a604fc..3553097e11 100644 --- a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[2]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RO { /* 0x0000 */ int arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) { diff --git a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl index 71c7a837af..97d28572b0 100644 --- a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RW { /* 0x0000 */ int arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) { diff --git a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl index ad28d92e06..ed949991fa 100644 --- a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[2]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RO { /* 0x0000 */ float arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) { diff --git a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl index fdbb89de82..525920c8df 100644 --- a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RW { /* 0x0000 */ float arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) { diff --git a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl index 14068e5729..97cbb4f133 100644 --- a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[2]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RO { /* 0x0000 */ uint arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) { diff --git a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl index 04f85089e2..d345929025 100644 --- a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl +++ b/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl @@ -2,7 +2,7 @@ using namespace metal; struct tint_symbol_2 { - /* 0x0000 */ uint buffer_size[1]; + /* 0x0000 */ uint4 buffer_size[1]; }; struct SB_RW { /* 0x0000 */ uint arg_0[1]; @@ -12,7 +12,7 @@ struct tint_symbol { }; void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) { - uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u); + uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u); } vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {