mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-12-21 02:39:11 +00:00
Remove support for the @stride attribute
We still use the StrideAttribute AST node in the SPIR-V reader for strided arrays and matrices, which are then removed by transforms. The WGSL parser no longer has to handle attributes on types. Bug: tint:1381 Change-Id: Ifa39575ce207d3fdfcbef7125fe6a3686fad5f20 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/83963 Reviewed-by: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
James Price
@@ -177,94 +177,6 @@ fn main() {
|
||||
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
|
||||
}
|
||||
|
||||
TEST_F(ArrayLengthFromUniformTest, WithStride) {
|
||||
auto* src = R"(
|
||||
@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = arrayLength(&sb);
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
struct tint_symbol {
|
||||
buffer_size : array<vec4<u32>, 1u>;
|
||||
}
|
||||
|
||||
@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 64u);
|
||||
}
|
||||
)";
|
||||
|
||||
ArrayLengthFromUniform::Config cfg({0, 30u});
|
||||
cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
|
||||
|
||||
DataMap data;
|
||||
data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
|
||||
|
||||
auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
EXPECT_EQ(std::unordered_set<uint32_t>({0}),
|
||||
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
|
||||
}
|
||||
|
||||
TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) {
|
||||
auto* src = R"(
|
||||
struct SB {
|
||||
x : i32;
|
||||
y : f32;
|
||||
arr : @stride(64) array<i32>;
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : SB;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = arrayLength(&sb.arr);
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
struct tint_symbol {
|
||||
buffer_size : array<vec4<u32>, 1u>;
|
||||
}
|
||||
|
||||
@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
|
||||
|
||||
struct SB {
|
||||
x : i32;
|
||||
y : f32;
|
||||
arr : @stride(64) array<i32>;
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : SB;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
|
||||
}
|
||||
)";
|
||||
|
||||
ArrayLengthFromUniform::Config cfg({0, 30u});
|
||||
cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
|
||||
|
||||
DataMap data;
|
||||
data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
|
||||
|
||||
auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
EXPECT_EQ(std::unordered_set<uint32_t>({0}),
|
||||
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
|
||||
}
|
||||
|
||||
TEST_F(ArrayLengthFromUniformTest, MultipleStorageBuffers) {
|
||||
auto* src = R"(
|
||||
struct SB1 {
|
||||
|
||||
@@ -287,78 +287,6 @@ fn main() {
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(CalculateArrayLengthTest, WithStride) {
|
||||
auto* src = R"(
|
||||
@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = arrayLength(&sb);
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
@internal(intrinsic_buffer_size)
|
||||
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : @stride(64) array<i32>, result : ptr<function, u32>)
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var tint_symbol_1 : u32 = 0u;
|
||||
tint_symbol(sb, &(tint_symbol_1));
|
||||
let tint_symbol_2 : u32 = (tint_symbol_1 / 64u);
|
||||
var len : u32 = tint_symbol_2;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(CalculateArrayLengthTest, WithStride_InStruct) {
|
||||
auto* src = R"(
|
||||
struct SB {
|
||||
x : i32;
|
||||
y : f32;
|
||||
arr : @stride(64) array<i32>;
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : SB;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var len : u32 = arrayLength(&sb.arr);
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
@internal(intrinsic_buffer_size)
|
||||
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
|
||||
|
||||
struct SB {
|
||||
x : i32;
|
||||
y : f32;
|
||||
arr : @stride(64) array<i32>;
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read> sb : SB;
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var tint_symbol_1 : u32 = 0u;
|
||||
tint_symbol(sb, &(tint_symbol_1));
|
||||
let tint_symbol_2 : u32 = ((tint_symbol_1 - 8u) / 64u);
|
||||
var len : u32 = tint_symbol_2;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(CalculateArrayLengthTest, Nested) {
|
||||
auto* src = R"(
|
||||
struct SB {
|
||||
|
||||
@@ -1970,22 +1970,26 @@ struct SB {
|
||||
|
||||
TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain) {
|
||||
auto* src = R"(
|
||||
// sizeof(S1) == 32
|
||||
// alignof(S1) == 16
|
||||
struct S1 {
|
||||
a : i32;
|
||||
b : vec3<f32>;
|
||||
c : i32;
|
||||
};
|
||||
|
||||
// sizeof(S2) == 116
|
||||
// alignof(S2) == 16
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
};
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -1999,9 +2003,9 @@ fn main() {
|
||||
// sb.b[4].b[1].b.z
|
||||
// ^ ^ ^ ^ ^ ^
|
||||
// | | | | | |
|
||||
// 128 | |1200| 1224
|
||||
// 128 | |688 | 712
|
||||
// | | |
|
||||
// 1152 1168 1216
|
||||
// 640 656 704
|
||||
|
||||
auto* expect = R"(
|
||||
struct S1 {
|
||||
@@ -2012,14 +2016,14 @@ struct S1 {
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
}
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2029,7 +2033,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var x : f32 = tint_symbol(sb, 1224u);
|
||||
var x : f32 = tint_symbol(sb, 712u);
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -2050,12 +2054,12 @@ fn main() {
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
};
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
};
|
||||
|
||||
@@ -2069,9 +2073,9 @@ struct S1 {
|
||||
// sb.b[4].b[1].b.z
|
||||
// ^ ^ ^ ^ ^ ^
|
||||
// | | | | | |
|
||||
// 128 | |1200| 1224
|
||||
// 128 | |688 | 712
|
||||
// | | |
|
||||
// 1152 1168 1216
|
||||
// 640 656 704
|
||||
|
||||
auto* expect = R"(
|
||||
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
|
||||
@@ -2079,7 +2083,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
|
||||
|
||||
@stage(compute) @workgroup_size(1)
|
||||
fn main() {
|
||||
var x : f32 = tint_symbol(sb, 1224u);
|
||||
var x : f32 = tint_symbol(sb, 712u);
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2087,12 +2091,12 @@ fn main() {
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
}
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
}
|
||||
|
||||
@@ -2118,14 +2122,14 @@ struct S1 {
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
};
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
};
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2148,14 +2152,14 @@ struct S1 {
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
}
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2168,7 +2172,7 @@ fn main() {
|
||||
var i : i32 = 4;
|
||||
var j : u32 = 1u;
|
||||
var k : i32 = 2;
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -2192,12 +2196,12 @@ fn main() {
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b :array<S2>;
|
||||
};
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
};
|
||||
|
||||
@@ -2217,7 +2221,7 @@ fn main() {
|
||||
var i : i32 = 4;
|
||||
var j : u32 = 1u;
|
||||
var k : i32 = 2;
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2225,12 +2229,12 @@ fn main() {
|
||||
struct SB {
|
||||
@size(128)
|
||||
a : i32;
|
||||
b : @stride(256) array<S2>;
|
||||
b : array<S2>;
|
||||
}
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
b : @stride(32) array<S1, 3>;
|
||||
b : array<S1, 3>;
|
||||
c : i32;
|
||||
}
|
||||
|
||||
@@ -2256,7 +2260,7 @@ struct S1 {
|
||||
|
||||
type A1 = S1;
|
||||
|
||||
type A1_Array = @stride(32) array<S1, 3>;
|
||||
type A1_Array = array<S1, 3>;
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
@@ -2266,7 +2270,7 @@ struct S2 {
|
||||
|
||||
type A2 = S2;
|
||||
|
||||
type A2_Array = @stride(256) array<S2>;
|
||||
type A2_Array = array<S2>;
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
@@ -2294,7 +2298,7 @@ struct S1 {
|
||||
|
||||
type A1 = S1;
|
||||
|
||||
type A1_Array = @stride(32) array<S1, 3>;
|
||||
type A1_Array = array<S1, 3>;
|
||||
|
||||
struct S2 {
|
||||
a : i32;
|
||||
@@ -2304,7 +2308,7 @@ struct S2 {
|
||||
|
||||
type A2 = S2;
|
||||
|
||||
type A2_Array = @stride(256) array<S2>;
|
||||
type A2_Array = array<S2>;
|
||||
|
||||
struct SB {
|
||||
@size(128)
|
||||
@@ -2322,7 +2326,7 @@ fn main() {
|
||||
var i : i32 = 4;
|
||||
var j : u32 = 1u;
|
||||
var k : i32 = 2;
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -2350,7 +2354,7 @@ struct SB {
|
||||
b : A2_Array;
|
||||
};
|
||||
|
||||
type A2_Array = @stride(256) array<S2>;
|
||||
type A2_Array = array<S2>;
|
||||
|
||||
type A2 = S2;
|
||||
|
||||
@@ -2362,7 +2366,7 @@ struct S2 {
|
||||
|
||||
type A1 = S1;
|
||||
|
||||
type A1_Array = @stride(32) array<S1, 3>;
|
||||
type A1_Array = array<S1, 3>;
|
||||
|
||||
struct S1 {
|
||||
a : i32;
|
||||
@@ -2380,7 +2384,7 @@ fn main() {
|
||||
var i : i32 = 4;
|
||||
var j : u32 = 1u;
|
||||
var k : i32 = 2;
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> sb : SB;
|
||||
@@ -2391,7 +2395,7 @@ struct SB {
|
||||
b : A2_Array;
|
||||
}
|
||||
|
||||
type A2_Array = @stride(256) array<S2>;
|
||||
type A2_Array = array<S2>;
|
||||
|
||||
type A2 = S2;
|
||||
|
||||
@@ -2403,7 +2407,7 @@ struct S2 {
|
||||
|
||||
type A1 = S1;
|
||||
|
||||
type A1_Array = @stride(32) array<S1, 3>;
|
||||
type A1_Array = array<S1, 3>;
|
||||
|
||||
struct S1 {
|
||||
a : i32;
|
||||
|
||||
@@ -27,7 +27,6 @@
|
||||
#include "src/tint/transform/fold_trivial_single_use_lets.h"
|
||||
#include "src/tint/transform/loop_to_for_loop.h"
|
||||
#include "src/tint/transform/manager.h"
|
||||
#include "src/tint/transform/pad_array_elements.h"
|
||||
#include "src/tint/transform/promote_initializers_to_const_var.h"
|
||||
#include "src/tint/transform/promote_side_effects_to_decl.h"
|
||||
#include "src/tint/transform/remove_phonies.h"
|
||||
@@ -105,7 +104,6 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) const {
|
||||
}
|
||||
manager.Add<PromoteInitializersToConstVar>();
|
||||
|
||||
manager.Add<PadArrayElements>();
|
||||
manager.Add<AddEmptyEntryPoint>();
|
||||
manager.Add<AddSpirvBlockAttribute>();
|
||||
|
||||
|
||||
@@ -1,177 +0,0 @@
|
||||
// Copyright 2021 The Tint Authors.
|
||||
//
|
||||
// Licensed under the Apache License, Version 2.0 (the "License");
|
||||
// you may not use this file except in compliance with the License.
|
||||
// You may obtain a copy of the License at
|
||||
//
|
||||
// http://www.apache.org/licenses/LICENSE-2.0
|
||||
//
|
||||
// Unless required by applicable law or agreed to in writing, software
|
||||
// distributed under the License is distributed on an "AS IS" BASIS,
|
||||
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/tint/transform/pad_array_elements.h"
|
||||
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
|
||||
#include "src/tint/program_builder.h"
|
||||
#include "src/tint/sem/array.h"
|
||||
#include "src/tint/sem/call.h"
|
||||
#include "src/tint/sem/expression.h"
|
||||
#include "src/tint/sem/type_constructor.h"
|
||||
#include "src/tint/utils/map.h"
|
||||
|
||||
TINT_INSTANTIATE_TYPEINFO(tint::transform::PadArrayElements);
|
||||
|
||||
namespace tint {
|
||||
namespace transform {
|
||||
namespace {
|
||||
|
||||
using ArrayBuilder = std::function<const ast::Array*()>;
|
||||
|
||||
/// PadArray returns a function that constructs a new array in `ctx.dst` with
|
||||
/// the element type padded to account for the explicit stride. PadArray will
|
||||
/// recursively pad arrays-of-arrays. The new array element type will be added
|
||||
/// to module-scope type declarations of `ctx.dst`.
|
||||
/// @param ctx the CloneContext
|
||||
/// @param create_ast_type_for Transform::CreateASTTypeFor()
|
||||
/// @param padded_arrays a map of src array type to the new array name
|
||||
/// @param array the array type
|
||||
/// @return the new AST array
|
||||
template <typename CREATE_AST_TYPE_FOR>
|
||||
ArrayBuilder PadArray(
|
||||
CloneContext& ctx,
|
||||
CREATE_AST_TYPE_FOR&& create_ast_type_for,
|
||||
std::unordered_map<const sem::Array*, ArrayBuilder>& padded_arrays,
|
||||
const sem::Array* array) {
|
||||
if (array->IsStrideImplicit()) {
|
||||
// We don't want to wrap arrays that have an implicit stride
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
return utils::GetOrCreate(padded_arrays, array, [&] {
|
||||
// Generate a unique name for the array element type
|
||||
auto name = ctx.dst->Symbols().New("tint_padded_array_element");
|
||||
|
||||
// Examine the element type. Is it also an array?
|
||||
const ast::Type* el_ty = nullptr;
|
||||
if (auto* el_array = array->ElemType()->As<sem::Array>()) {
|
||||
// Array of array - call PadArray() on the element type
|
||||
if (auto p =
|
||||
PadArray(ctx, create_ast_type_for, padded_arrays, el_array)) {
|
||||
el_ty = p();
|
||||
}
|
||||
}
|
||||
|
||||
// If the element wasn't a padded array, just create the typical AST type
|
||||
// for it
|
||||
if (el_ty == nullptr) {
|
||||
el_ty = create_ast_type_for(ctx, array->ElemType());
|
||||
}
|
||||
|
||||
// Structure() will create and append the ast::Struct to the
|
||||
// global declarations of `ctx.dst`. As we haven't finished building the
|
||||
// current module-scope statement or function, this will be placed
|
||||
// immediately before the usage.
|
||||
ctx.dst->Structure(
|
||||
name,
|
||||
{ctx.dst->Member("el", el_ty, {ctx.dst->MemberSize(array->Stride())})});
|
||||
|
||||
auto* dst = ctx.dst;
|
||||
return [=] {
|
||||
if (array->IsRuntimeSized()) {
|
||||
return dst->ty.array(dst->create<ast::TypeName>(name));
|
||||
} else {
|
||||
return dst->ty.array(dst->create<ast::TypeName>(name), array->Count());
|
||||
}
|
||||
};
|
||||
});
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
PadArrayElements::PadArrayElements() = default;
|
||||
|
||||
PadArrayElements::~PadArrayElements() = default;
|
||||
|
||||
bool PadArrayElements::ShouldRun(const Program* program, const DataMap&) const {
|
||||
for (auto* node : program->ASTNodes().Objects()) {
|
||||
if (auto* var = node->As<ast::Type>()) {
|
||||
if (auto* arr = program->Sem().Get<sem::Array>(var)) {
|
||||
if (!arr->IsStrideImplicit()) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
void PadArrayElements::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
|
||||
auto& sem = ctx.src->Sem();
|
||||
|
||||
std::unordered_map<const sem::Array*, ArrayBuilder> padded_arrays;
|
||||
auto pad = [&](const sem::Array* array) {
|
||||
return PadArray(ctx, CreateASTTypeFor, padded_arrays, array);
|
||||
};
|
||||
|
||||
// Replace all array types with their corresponding padded array type
|
||||
ctx.ReplaceAll([&](const ast::Type* ast_type) -> const ast::Type* {
|
||||
auto* type = ctx.src->TypeOf(ast_type);
|
||||
if (auto* array = type->UnwrapRef()->As<sem::Array>()) {
|
||||
if (auto p = pad(array)) {
|
||||
return p();
|
||||
}
|
||||
}
|
||||
return nullptr;
|
||||
});
|
||||
|
||||
// Fix up index accessors so `a[1]` becomes `a[1].el`
|
||||
ctx.ReplaceAll([&](const ast::IndexAccessorExpression* accessor)
|
||||
-> const ast::Expression* {
|
||||
if (auto* array = tint::As<sem::Array>(
|
||||
sem.Get(accessor->object)->Type()->UnwrapRef())) {
|
||||
if (pad(array)) {
|
||||
// Array element is wrapped in a structure. Emit a member accessor
|
||||
// to get to the actual array element.
|
||||
auto* idx = ctx.CloneWithoutTransform(accessor);
|
||||
return ctx.dst->MemberAccessor(idx, "el");
|
||||
}
|
||||
}
|
||||
return nullptr;
|
||||
});
|
||||
|
||||
// Fix up array constructors so `A(1,2)` becomes
|
||||
// `A(padded(1), padded(2))`
|
||||
ctx.ReplaceAll(
|
||||
[&](const ast::CallExpression* expr) -> const ast::Expression* {
|
||||
auto* call = sem.Get(expr);
|
||||
if (auto* ctor = call->Target()->As<sem::TypeConstructor>()) {
|
||||
if (auto* array = ctor->ReturnType()->As<sem::Array>()) {
|
||||
if (auto p = pad(array)) {
|
||||
auto* arr_ty = p();
|
||||
auto el_typename = arr_ty->type->As<ast::TypeName>()->name;
|
||||
|
||||
ast::ExpressionList args;
|
||||
args.reserve(call->Arguments().size());
|
||||
for (auto* arg : call->Arguments()) {
|
||||
auto* val = ctx.Clone(arg->Declaration());
|
||||
args.emplace_back(ctx.dst->Construct(
|
||||
ctx.dst->create<ast::TypeName>(el_typename), val));
|
||||
}
|
||||
|
||||
return ctx.dst->Construct(arr_ty, args);
|
||||
}
|
||||
}
|
||||
}
|
||||
return nullptr;
|
||||
});
|
||||
|
||||
ctx.Clone();
|
||||
}
|
||||
|
||||
} // namespace transform
|
||||
} // namespace tint
|
||||
@@ -1,62 +0,0 @@
|
||||
// Copyright 2021 The Tint Authors.
|
||||
//
|
||||
// Licensed under the Apache License, Version 2.0 (the "License");
|
||||
// you may not use this file except in compliance with the License.
|
||||
// You may obtain a copy of the License at
|
||||
//
|
||||
// http://www.apache.org/licenses/LICENSE-2.0
|
||||
//
|
||||
// Unless required by applicable law or agreed to in writing, software
|
||||
// distributed under the License is distributed on an "AS IS" BASIS,
|
||||
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#ifndef SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
|
||||
#define SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
|
||||
|
||||
#include "src/tint/transform/transform.h"
|
||||
|
||||
namespace tint {
|
||||
namespace transform {
|
||||
|
||||
/// PadArrayElements is a transform that replaces array types with an explicit
|
||||
/// stride that is larger than the implicit stride, with an array of a new
|
||||
/// structure type. This structure holds with a single field of the element
|
||||
/// type, decorated with a `@size` attribute to pad the structure to the
|
||||
/// required array stride. The new array types have no explicit stride,
|
||||
/// structure size is equal to the desired stride.
|
||||
/// Array index expressions and constructors are also adjusted to deal with this
|
||||
/// structure element type.
|
||||
/// This transform helps with backends that cannot directly return arrays or use
|
||||
/// them as parameters.
|
||||
class PadArrayElements : public Castable<PadArrayElements, Transform> {
|
||||
public:
|
||||
/// Constructor
|
||||
PadArrayElements();
|
||||
|
||||
/// Destructor
|
||||
~PadArrayElements() override;
|
||||
|
||||
/// @param program the program to inspect
|
||||
/// @param data optional extra transform-specific input data
|
||||
/// @returns true if this transform should be run for the given program
|
||||
bool ShouldRun(const Program* program,
|
||||
const DataMap& data = {}) const override;
|
||||
|
||||
protected:
|
||||
/// Runs the transform using the CloneContext built for transforming a
|
||||
/// program. Run() is responsible for calling Clone() on the CloneContext.
|
||||
/// @param ctx the CloneContext primed with the input program and
|
||||
/// ProgramBuilder
|
||||
/// @param inputs optional extra transform-specific input data
|
||||
/// @param outputs optional extra transform-specific output data
|
||||
void Run(CloneContext& ctx,
|
||||
const DataMap& inputs,
|
||||
DataMap& outputs) const override;
|
||||
};
|
||||
|
||||
} // namespace transform
|
||||
} // namespace tint
|
||||
|
||||
#endif // SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
|
||||
@@ -1,518 +0,0 @@
|
||||
// Copyright 2021 The Tint Authors.
|
||||
//
|
||||
// Licensed under the Apache License, Version 2.0 (the "License");
|
||||
// you may not use this file except in compliance with the License.
|
||||
// You may obtain a copy of the License at
|
||||
//
|
||||
// http://www.apache.org/licenses/LICENSE-2.0
|
||||
//
|
||||
// Unless required by applicable law or agreed to in writing, software
|
||||
// distributed under the License is distributed on an "AS IS" BASIS,
|
||||
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
// See the License for the specific language governing permissions and
|
||||
// limitations under the License.
|
||||
|
||||
#include "src/tint/transform/pad_array_elements.h"
|
||||
|
||||
#include <utility>
|
||||
|
||||
#include "src/tint/transform/test_helper.h"
|
||||
|
||||
namespace tint {
|
||||
namespace transform {
|
||||
namespace {
|
||||
|
||||
using PadArrayElementsTest = TransformTest;
|
||||
|
||||
TEST_F(PadArrayElementsTest, ShouldRunEmptyModule) {
|
||||
auto* src = R"()";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ShouldRunHasImplicitArrayStride) {
|
||||
auto* src = R"(
|
||||
var<private> arr : array<i32, 4>;
|
||||
)";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ShouldRunHasExplicitArrayStride) {
|
||||
auto* src = R"(
|
||||
var<private> arr : @stride(8) array<i32, 4>;
|
||||
)";
|
||||
|
||||
EXPECT_TRUE(ShouldRun<PadArrayElements>(src));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, EmptyModule) {
|
||||
auto* src = "";
|
||||
auto* expect = "";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ImplicitArrayStride) {
|
||||
auto* src = R"(
|
||||
var<private> arr : array<i32, 4>;
|
||||
)";
|
||||
auto* expect = src;
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArrayAsGlobal) {
|
||||
auto* src = R"(
|
||||
var<private> arr : @stride(8) array<i32, 4>;
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
var<private> arr : array<tint_padded_array_element, 4u>;
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, RuntimeArray) {
|
||||
auto* src = R"(
|
||||
struct S {
|
||||
rta : @stride(8) array<i32>;
|
||||
};
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct S {
|
||||
rta : array<tint_padded_array_element>;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArrayFunctionVar) {
|
||||
auto* src = R"(
|
||||
fn f() {
|
||||
var arr : @stride(16) array<i32, 4>;
|
||||
arr = @stride(16) array<i32, 4>();
|
||||
arr = @stride(16) array<i32, 4>(1, 2, 3, 4);
|
||||
let x = arr[3];
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(16)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
fn f() {
|
||||
var arr : array<tint_padded_array_element, 4u>;
|
||||
arr = array<tint_padded_array_element, 4u>();
|
||||
arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
|
||||
let x = arr[3].el;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArrayAsParam) {
|
||||
auto* src = R"(
|
||||
fn f(a : @stride(12) array<i32, 4>) -> i32 {
|
||||
return a[2];
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(12)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
fn f(a : array<tint_padded_array_element, 4u>) -> i32 {
|
||||
return a[2].el;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
// TODO(crbug.com/tint/781): Cannot parse the stride on the return array type.
|
||||
TEST_F(PadArrayElementsTest, DISABLED_ArrayAsReturn) {
|
||||
auto* src = R"(
|
||||
fn f() -> @stride(8) array<i32, 4> {
|
||||
return array<i32, 4>(1, 2, 3, 4);
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
el : i32;
|
||||
@size(4)
|
||||
padding : u32;
|
||||
};
|
||||
|
||||
fn f() -> array<tint_padded_array_element, 4> {
|
||||
return array<tint_padded_array_element, 4>(tint_padded_array_element(1, 0u), tint_padded_array_element(2, 0u), tint_padded_array_element(3, 0u), tint_padded_array_element(4, 0u));
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArrayAlias) {
|
||||
auto* src = R"(
|
||||
type Array = @stride(16) array<i32, 4>;
|
||||
|
||||
fn f() {
|
||||
var arr : Array;
|
||||
arr = Array();
|
||||
arr = Array(1, 2, 3, 4);
|
||||
let vals : Array = Array(1, 2, 3, 4);
|
||||
arr = vals;
|
||||
let x = arr[3];
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(16)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
type Array = array<tint_padded_array_element, 4u>;
|
||||
|
||||
fn f() {
|
||||
var arr : array<tint_padded_array_element, 4u>;
|
||||
arr = array<tint_padded_array_element, 4u>();
|
||||
arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
|
||||
let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
|
||||
arr = vals;
|
||||
let x = arr[3].el;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArrayAlias_OutOfOrder) {
|
||||
auto* src = R"(
|
||||
fn f() {
|
||||
var arr : Array;
|
||||
arr = Array();
|
||||
arr = Array(1, 2, 3, 4);
|
||||
let vals : Array = Array(1, 2, 3, 4);
|
||||
arr = vals;
|
||||
let x = arr[3];
|
||||
}
|
||||
|
||||
type Array = @stride(16) array<i32, 4>;
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(16)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
fn f() {
|
||||
var arr : array<tint_padded_array_element, 4u>;
|
||||
arr = array<tint_padded_array_element, 4u>();
|
||||
arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
|
||||
let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
|
||||
arr = vals;
|
||||
let x = arr[3].el;
|
||||
}
|
||||
|
||||
type Array = array<tint_padded_array_element, 4u>;
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArraysInStruct) {
|
||||
auto* src = R"(
|
||||
struct S {
|
||||
a : @stride(8) array<i32, 4>;
|
||||
b : @stride(8) array<i32, 8>;
|
||||
c : @stride(8) array<i32, 4>;
|
||||
d : @stride(12) array<i32, 8>;
|
||||
};
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_1 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_2 {
|
||||
@size(12)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct S {
|
||||
a : array<tint_padded_array_element, 4u>;
|
||||
b : array<tint_padded_array_element_1, 8u>;
|
||||
c : array<tint_padded_array_element, 4u>;
|
||||
d : array<tint_padded_array_element_2, 8u>;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, ArraysOfArraysInStruct) {
|
||||
auto* src = R"(
|
||||
struct S {
|
||||
a : @stride(512) array<i32, 4>;
|
||||
b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
|
||||
c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
|
||||
};
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(512)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_2 {
|
||||
@size(32)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_1 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_2, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_5 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_4 {
|
||||
@size(64)
|
||||
el : array<tint_padded_array_element_5, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_3 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_4, 4u>;
|
||||
}
|
||||
|
||||
struct S {
|
||||
a : array<tint_padded_array_element, 4u>;
|
||||
b : array<tint_padded_array_element_1, 4u>;
|
||||
c : array<tint_padded_array_element_3, 4u>;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct) {
|
||||
auto* src = R"(
|
||||
struct S {
|
||||
a : @stride(512) array<i32, 4>;
|
||||
b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
|
||||
c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
|
||||
};
|
||||
|
||||
fn f(s : S) -> i32 {
|
||||
return s.a[2] + s.b[1][2] + s.c[3][1][2];
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(512)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_2 {
|
||||
@size(32)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_1 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_2, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_5 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_4 {
|
||||
@size(64)
|
||||
el : array<tint_padded_array_element_5, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_3 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_4, 4u>;
|
||||
}
|
||||
|
||||
struct S {
|
||||
a : array<tint_padded_array_element, 4u>;
|
||||
b : array<tint_padded_array_element_1, 4u>;
|
||||
c : array<tint_padded_array_element_3, 4u>;
|
||||
}
|
||||
|
||||
fn f(s : S) -> i32 {
|
||||
return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct_OutOfOrder) {
|
||||
auto* src = R"(
|
||||
fn f(s : S) -> i32 {
|
||||
return s.a[2] + s.b[1][2] + s.c[3][1][2];
|
||||
}
|
||||
|
||||
struct S {
|
||||
a : @stride(512) array<i32, 4>;
|
||||
b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
|
||||
c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
|
||||
};
|
||||
)";
|
||||
auto* expect = R"(
|
||||
struct tint_padded_array_element {
|
||||
@size(512)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_1 {
|
||||
@size(32)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_2 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_1, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_3 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_4 {
|
||||
@size(64)
|
||||
el : array<tint_padded_array_element_3, 4u>;
|
||||
}
|
||||
|
||||
struct tint_padded_array_element_5 {
|
||||
@size(512)
|
||||
el : array<tint_padded_array_element_4, 4u>;
|
||||
}
|
||||
|
||||
fn f(s : S) -> i32 {
|
||||
return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
|
||||
}
|
||||
|
||||
struct S {
|
||||
a : array<tint_padded_array_element, 4u>;
|
||||
b : array<tint_padded_array_element_2, 4u>;
|
||||
c : array<tint_padded_array_element_5, 4u>;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(PadArrayElementsTest, DeclarationOrder) {
|
||||
auto* src = R"(
|
||||
type T0 = i32;
|
||||
|
||||
type T1 = @stride(8) array<i32, 1>;
|
||||
|
||||
type T2 = i32;
|
||||
|
||||
fn f1(a : @stride(8) array<i32, 2>) {
|
||||
}
|
||||
|
||||
type T3 = i32;
|
||||
|
||||
fn f2() {
|
||||
var v : @stride(8) array<i32, 3>;
|
||||
}
|
||||
)";
|
||||
auto* expect = R"(
|
||||
type T0 = i32;
|
||||
|
||||
struct tint_padded_array_element {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
type T1 = array<tint_padded_array_element, 1u>;
|
||||
|
||||
type T2 = i32;
|
||||
|
||||
struct tint_padded_array_element_1 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
fn f1(a : array<tint_padded_array_element_1, 2u>) {
|
||||
}
|
||||
|
||||
type T3 = i32;
|
||||
|
||||
struct tint_padded_array_element_2 {
|
||||
@size(8)
|
||||
el : i32;
|
||||
}
|
||||
|
||||
fn f2() {
|
||||
var v : array<tint_padded_array_element_2, 3u>;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<PadArrayElements>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
} // namespace
|
||||
} // namespace transform
|
||||
} // namespace tint
|
||||
@@ -1431,7 +1431,7 @@ struct S {
|
||||
};
|
||||
@group(0) @binding(0) var<storage, read> s : S;
|
||||
|
||||
type UArr = @stride(16) array<f32, 4>;
|
||||
type UArr = array<vec4<f32>, 4>;
|
||||
struct U {
|
||||
a : UArr;
|
||||
};
|
||||
@@ -1451,11 +1451,11 @@ fn f() {
|
||||
var i32_sb4 : f32 = s.b[-1];
|
||||
var i32_sb5 : f32 = s.b[-4];
|
||||
|
||||
var i32_ua1 : f32 = u.a[4];
|
||||
var i32_ua2 : f32 = u.a[1];
|
||||
var i32_ua3 : f32 = u.a[0];
|
||||
var i32_ua4 : f32 = u.a[-1];
|
||||
var i32_ua5 : f32 = u.a[-4];
|
||||
var i32_ua1 : f32 = u.a[4].x;
|
||||
var i32_ua2 : f32 = u.a[1].x;
|
||||
var i32_ua3 : f32 = u.a[0].x;
|
||||
var i32_ua4 : f32 = u.a[-1].x;
|
||||
var i32_ua5 : f32 = u.a[-4].x;
|
||||
|
||||
// Unsigned
|
||||
var u32_sa1 : f32 = s.a[0u];
|
||||
@@ -1472,12 +1472,12 @@ fn f() {
|
||||
var u32_sb5 : f32 = s.b[10u];
|
||||
var u32_sb6 : f32 = s.b[100u];
|
||||
|
||||
var u32_ua1 : f32 = u.a[0u];
|
||||
var u32_ua2 : f32 = u.a[1u];
|
||||
var u32_ua3 : f32 = u.a[3u];
|
||||
var u32_ua4 : f32 = u.a[4u];
|
||||
var u32_ua5 : f32 = u.a[10u];
|
||||
var u32_ua6 : f32 = u.a[100u];
|
||||
var u32_ua1 : f32 = u.a[0u].x;
|
||||
var u32_ua2 : f32 = u.a[1u].x;
|
||||
var u32_ua3 : f32 = u.a[3u].x;
|
||||
var u32_ua4 : f32 = u.a[4u].x;
|
||||
var u32_ua5 : f32 = u.a[10u].x;
|
||||
var u32_ua6 : f32 = u.a[100u].x;
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -1490,7 +1490,7 @@ struct S {
|
||||
|
||||
@group(0) @binding(0) var<storage, read> s : S;
|
||||
|
||||
type UArr = @stride(16) array<f32, 4>;
|
||||
type UArr = array<vec4<f32>, 4>;
|
||||
|
||||
struct U {
|
||||
a : UArr;
|
||||
@@ -1509,11 +1509,11 @@ fn f() {
|
||||
var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_ua1 : f32 = u.a[3];
|
||||
var i32_ua2 : f32 = u.a[1];
|
||||
var i32_ua3 : f32 = u.a[0];
|
||||
var i32_ua4 : f32 = u.a[0];
|
||||
var i32_ua5 : f32 = u.a[0];
|
||||
var i32_ua1 : f32 = u.a[3].x;
|
||||
var i32_ua2 : f32 = u.a[1].x;
|
||||
var i32_ua3 : f32 = u.a[0].x;
|
||||
var i32_ua4 : f32 = u.a[0].x;
|
||||
var i32_ua5 : f32 = u.a[0].x;
|
||||
var u32_sa1 : f32 = s.a[0u];
|
||||
var u32_sa2 : f32 = s.a[1u];
|
||||
var u32_sa3 : f32 = s.a[3u];
|
||||
@@ -1526,12 +1526,12 @@ fn f() {
|
||||
var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_ua1 : f32 = u.a[0u];
|
||||
var u32_ua2 : f32 = u.a[1u];
|
||||
var u32_ua3 : f32 = u.a[3u];
|
||||
var u32_ua4 : f32 = u.a[3u];
|
||||
var u32_ua5 : f32 = u.a[3u];
|
||||
var u32_ua6 : f32 = u.a[3u];
|
||||
var u32_ua1 : f32 = u.a[0u].x;
|
||||
var u32_ua2 : f32 = u.a[1u].x;
|
||||
var u32_ua3 : f32 = u.a[3u].x;
|
||||
var u32_ua4 : f32 = u.a[3u].x;
|
||||
var u32_ua5 : f32 = u.a[3u].x;
|
||||
var u32_ua6 : f32 = u.a[3u].x;
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -1553,7 +1553,7 @@ struct S {
|
||||
|
||||
@group(0) @binding(0) var<storage, read> s : S;
|
||||
|
||||
type UArr = @stride(16) array<f32, 4>;
|
||||
type UArr = array<vec4<f32>, 4>;
|
||||
|
||||
struct U {
|
||||
a : UArr;
|
||||
@@ -1572,11 +1572,11 @@ fn f() {
|
||||
var i32_sb3 : f32 = s.b[0];
|
||||
var i32_sb4 : f32 = s.b[-1];
|
||||
var i32_sb5 : f32 = s.b[-4];
|
||||
var i32_ua1 : f32 = u.a[3];
|
||||
var i32_ua2 : f32 = u.a[1];
|
||||
var i32_ua3 : f32 = u.a[0];
|
||||
var i32_ua4 : f32 = u.a[0];
|
||||
var i32_ua5 : f32 = u.a[0];
|
||||
var i32_ua1 : f32 = u.a[3].x;
|
||||
var i32_ua2 : f32 = u.a[1].x;
|
||||
var i32_ua3 : f32 = u.a[0].x;
|
||||
var i32_ua4 : f32 = u.a[0].x;
|
||||
var i32_ua5 : f32 = u.a[0].x;
|
||||
var u32_sa1 : f32 = s.a[0u];
|
||||
var u32_sa2 : f32 = s.a[1u];
|
||||
var u32_sa3 : f32 = s.a[3u];
|
||||
@@ -1589,12 +1589,12 @@ fn f() {
|
||||
var u32_sb4 : f32 = s.b[4u];
|
||||
var u32_sb5 : f32 = s.b[10u];
|
||||
var u32_sb6 : f32 = s.b[100u];
|
||||
var u32_ua1 : f32 = u.a[0u];
|
||||
var u32_ua2 : f32 = u.a[1u];
|
||||
var u32_ua3 : f32 = u.a[3u];
|
||||
var u32_ua4 : f32 = u.a[3u];
|
||||
var u32_ua5 : f32 = u.a[3u];
|
||||
var u32_ua6 : f32 = u.a[3u];
|
||||
var u32_ua1 : f32 = u.a[0u].x;
|
||||
var u32_ua2 : f32 = u.a[1u].x;
|
||||
var u32_ua3 : f32 = u.a[3u].x;
|
||||
var u32_ua4 : f32 = u.a[3u].x;
|
||||
var u32_ua5 : f32 = u.a[3u].x;
|
||||
var u32_ua6 : f32 = u.a[3u].x;
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -1618,7 +1618,7 @@ struct S {
|
||||
|
||||
@group(0) @binding(0) var<storage, read> s : S;
|
||||
|
||||
type UArr = @stride(16) array<f32, 4>;
|
||||
type UArr = array<vec4<f32>, 4>;
|
||||
|
||||
struct U {
|
||||
a : UArr;
|
||||
@@ -1637,11 +1637,11 @@ fn f() {
|
||||
var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
|
||||
var i32_ua1 : f32 = u.a[4];
|
||||
var i32_ua2 : f32 = u.a[1];
|
||||
var i32_ua3 : f32 = u.a[0];
|
||||
var i32_ua4 : f32 = u.a[-1];
|
||||
var i32_ua5 : f32 = u.a[-4];
|
||||
var i32_ua1 : f32 = u.a[4].x;
|
||||
var i32_ua2 : f32 = u.a[1].x;
|
||||
var i32_ua3 : f32 = u.a[0].x;
|
||||
var i32_ua4 : f32 = u.a[-1].x;
|
||||
var i32_ua5 : f32 = u.a[-4].x;
|
||||
var u32_sa1 : f32 = s.a[0u];
|
||||
var u32_sa2 : f32 = s.a[1u];
|
||||
var u32_sa3 : f32 = s.a[3u];
|
||||
@@ -1654,12 +1654,12 @@ fn f() {
|
||||
var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
|
||||
var u32_ua1 : f32 = u.a[0u];
|
||||
var u32_ua2 : f32 = u.a[1u];
|
||||
var u32_ua3 : f32 = u.a[3u];
|
||||
var u32_ua4 : f32 = u.a[4u];
|
||||
var u32_ua5 : f32 = u.a[10u];
|
||||
var u32_ua6 : f32 = u.a[100u];
|
||||
var u32_ua1 : f32 = u.a[0u].x;
|
||||
var u32_ua2 : f32 = u.a[1u].x;
|
||||
var u32_ua3 : f32 = u.a[3u].x;
|
||||
var u32_ua4 : f32 = u.a[4u].x;
|
||||
var u32_ua5 : f32 = u.a[10u].x;
|
||||
var u32_ua6 : f32 = u.a[100u].x;
|
||||
}
|
||||
)";
|
||||
|
||||
@@ -1683,7 +1683,7 @@ struct S {
|
||||
|
||||
@group(0) @binding(0) var<storage, read> s : S;
|
||||
|
||||
type UArr = @stride(16) array<f32, 4>;
|
||||
type UArr = array<vec4<f32>, 4>;
|
||||
|
||||
struct U {
|
||||
a : UArr;
|
||||
@@ -1702,11 +1702,11 @@ fn f() {
|
||||
var i32_sb3 : f32 = s.b[0];
|
||||
var i32_sb4 : f32 = s.b[-1];
|
||||
var i32_sb5 : f32 = s.b[-4];
|
||||
var i32_ua1 : f32 = u.a[4];
|
||||
var i32_ua2 : f32 = u.a[1];
|
||||
var i32_ua3 : f32 = u.a[0];
|
||||
var i32_ua4 : f32 = u.a[-1];
|
||||
var i32_ua5 : f32 = u.a[-4];
|
||||
var i32_ua1 : f32 = u.a[4].x;
|
||||
var i32_ua2 : f32 = u.a[1].x;
|
||||
var i32_ua3 : f32 = u.a[0].x;
|
||||
var i32_ua4 : f32 = u.a[-1].x;
|
||||
var i32_ua5 : f32 = u.a[-4].x;
|
||||
var u32_sa1 : f32 = s.a[0u];
|
||||
var u32_sa2 : f32 = s.a[1u];
|
||||
var u32_sa3 : f32 = s.a[3u];
|
||||
@@ -1719,12 +1719,12 @@ fn f() {
|
||||
var u32_sb4 : f32 = s.b[4u];
|
||||
var u32_sb5 : f32 = s.b[10u];
|
||||
var u32_sb6 : f32 = s.b[100u];
|
||||
var u32_ua1 : f32 = u.a[0u];
|
||||
var u32_ua2 : f32 = u.a[1u];
|
||||
var u32_ua3 : f32 = u.a[3u];
|
||||
var u32_ua4 : f32 = u.a[4u];
|
||||
var u32_ua5 : f32 = u.a[10u];
|
||||
var u32_ua6 : f32 = u.a[100u];
|
||||
var u32_ua1 : f32 = u.a[0u].x;
|
||||
var u32_ua2 : f32 = u.a[1u].x;
|
||||
var u32_ua3 : f32 = u.a[3u].x;
|
||||
var u32_ua4 : f32 = u.a[4u].x;
|
||||
var u32_ua5 : f32 = u.a[10u].x;
|
||||
var u32_ua6 : f32 = u.a[100u].x;
|
||||
}
|
||||
)";
|
||||
|
||||
|
||||
@@ -258,7 +258,7 @@ struct State {
|
||||
ctx.dst->Symbols().New(kStructName),
|
||||
{
|
||||
ctx.dst->Member(GetStructBufferName(),
|
||||
ctx.dst->ty.array<ProgramBuilder::u32>(4)),
|
||||
ctx.dst->ty.array<ProgramBuilder::u32>()),
|
||||
});
|
||||
for (uint32_t i = 0; i < cfg.vertex_state.size(); ++i) {
|
||||
// The decorated variable with struct type
|
||||
|
||||
@@ -108,7 +108,7 @@ fn main() -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@stage(vertex)
|
||||
@@ -137,7 +137,7 @@ fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -175,7 +175,7 @@ fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -213,7 +213,7 @@ fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(5) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -256,7 +256,7 @@ fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -304,7 +304,7 @@ fn main(@location(0) var_a : f32,
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -364,7 +364,7 @@ fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -444,7 +444,7 @@ struct Inputs {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -526,7 +526,7 @@ fn main(inputs : Inputs, indices : Indices) -> @builtin(position) vec4<f32> {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -603,7 +603,7 @@ struct Indices {
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -670,7 +670,7 @@ fn main(@location(0) var_a : f32,
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -715,7 +715,7 @@ fn main(@location(0) var_a : vec2<f32>,
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -771,7 +771,7 @@ fn main(@location(0) var_a : f32,
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data_1 : @stride(4) array<u32>;
|
||||
tint_vertex_data_1 : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0_1 : TintVertexData;
|
||||
@@ -848,7 +848,7 @@ fn main(
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -1008,7 +1008,7 @@ fn main(
|
||||
auto* expect =
|
||||
R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
@@ -1167,7 +1167,7 @@ fn main(
|
||||
|
||||
auto* expect = R"(
|
||||
struct TintVertexData {
|
||||
tint_vertex_data : @stride(4) array<u32>;
|
||||
tint_vertex_data : array<u32>;
|
||||
}
|
||||
|
||||
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
|
||||
|
||||
Reference in New Issue
Block a user