encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Add transform/WrapArraysInStructs 

And replace the MSL writer's logic to do this with the transform.

We need to do the same thing in HLSL, and in the future GLSL too.

Partially reverts fbfde720

Change-Id: Ie280e011bc3ded8e15ccacc0aeb12da3c2407389
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/54242
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton 2021-06-16 09:19:36 +00:00 committed by Ben Clayton
parent cfed1cb01e
commit 0597a2b51b
31 changed files with 771 additions and 255 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/BUILD.gn
View File

@ -575,6 +575,8 @@ libtint_source_set("libtint_core_all_src") {
"transform/var_for_dynamic_index.h",
"transform/vertex_pulling.cc",
"transform/vertex_pulling.h",
"transform/wrap_arrays_in_structs.cc",
"transform/wrap_arrays_in_structs.h",
"utils/enum_set.h",
"utils/get_or_create.h",
"utils/hash.h",

3
src/CMakeLists.txt
View File

@ -300,6 +300,8 @@ set(TINT_LIB_SRCS
transform/var_for_dynamic_index.h
transform/vertex_pulling.cc
transform/vertex_pulling.h
transform/wrap_arrays_in_structs.cc
transform/wrap_arrays_in_structs.h
sem/bool_type.cc
sem/bool_type.h
sem/depth_texture_type.cc
@ -859,6 +861,7 @@ if(${TINT_BUILD_TESTS})
transform/test_helper.h
transform/var_for_dynamic_index_test.cc
transform/vertex_pulling_test.cc
transform/wrap_arrays_in_structs_test.cc
)
endif()

12
src/program_builder.h
View File

@ -1486,6 +1486,18 @@ class ProgramBuilder {
Expr(std::forward<RHS>(rhs)));
}
/// @param source the source information
/// @param arr the array argument for the array accessor expression
/// @param idx the index argument for the array accessor expression
/// @returns a `ast::ArrayAccessorExpression` that indexes `arr` with `idx`
template <typename ARR, typename IDX>
ast::ArrayAccessorExpression* IndexAccessor(const Source& source,
ARR&& arr,
IDX&& idx) {
return create<ast::ArrayAccessorExpression>(
source, Expr(std::forward<ARR>(arr)), Expr(std::forward<IDX>(idx)));
}
/// @param arr the array argument for the array accessor expression
/// @param idx the index argument for the array accessor expression
/// @returns a `ast::ArrayAccessorExpression` that indexes `arr` with `idx`

2
src/transform/msl.cc
View File

@ -28,6 +28,7 @@
#include "src/transform/external_texture_transform.h"
#include "src/transform/manager.h"
#include "src/transform/promote_initializers_to_const_var.h"
#include "src/transform/wrap_arrays_in_structs.h"
namespace tint {
namespace transform {
@ -41,6 +42,7 @@ Output Msl::Run(const Program* in, const DataMap&) {
manager.Add<CanonicalizeEntryPointIO>();
manager.Add<ExternalTextureTransform>();
manager.Add<PromoteInitializersToConstVar>();
manager.Add<WrapArraysInStructs>();
data.Add<CanonicalizeEntryPointIO::Config>(
CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
auto out = manager.Run(in, data);

152
src/transform/wrap_arrays_in_structs.cc Normal file
View File

@ -0,0 +1,152 @@
// 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/transform/wrap_arrays_in_structs.h"
#include <utility>
#include "src/program_builder.h"
#include "src/sem/array.h"
#include "src/sem/expression.h"
#include "src/utils/get_or_create.h"
namespace tint {
namespace transform {
WrapArraysInStructs::WrappedArrayInfo::WrappedArrayInfo() = default;
WrapArraysInStructs::WrappedArrayInfo::WrappedArrayInfo(
const WrappedArrayInfo&) = default;
WrapArraysInStructs::WrappedArrayInfo::~WrappedArrayInfo() = default;
WrapArraysInStructs::WrapArraysInStructs() = default;
WrapArraysInStructs::~WrapArraysInStructs() = default;
Output WrapArraysInStructs::Run(const Program* in, const DataMap&) {
ProgramBuilder out;
CloneContext ctx(&out, in);
auto& sem = ctx.src->Sem();
std::unordered_map<const sem::Array*, WrappedArrayInfo> wrapped_arrays;
auto wrapper = [&](const sem::Array* array) {
return WrapArray(ctx, wrapped_arrays, array);
};
auto wrapper_typename = [&](const sem::Array* arr) -> ast::TypeName* {
auto info = wrapper(arr);
return info ? ctx.dst->create<ast::TypeName>(info.wrapper_name) : nullptr;
};
// Replace all array types with their corresponding wrapper
ctx.ReplaceAll([&](ast::Type* ast_type) -> ast::Type* {
auto* type = ctx.src->TypeOf(ast_type);
if (auto* array = type->UnwrapRef()->As<sem::Array>()) {
return wrapper_typename(array);
}
return nullptr;
});
// Fix up array accessors so `a[1]` becomes `a.arr[1]`
ctx.ReplaceAll([&](ast::ArrayAccessorExpression* accessor)
-> ast::ArrayAccessorExpression* {
if (auto* array =
As<sem::Array>(sem.Get(accessor->array())->Type()->UnwrapRef())) {
if (wrapper(array)) {
// Array is wrapped in a structure. Emit a member accessor to get
// to the actual array.
auto* arr = ctx.Clone(accessor->array());
auto* idx = ctx.Clone(accessor->idx_expr());
auto* unwrapped = ctx.dst->MemberAccessor(arr, "arr");
return ctx.dst->IndexAccessor(accessor->source(), unwrapped, idx);
}
}
return nullptr;
});
// Fix up array constructors so `A(1,2)` becomes `tint_array_wrapper(A(1,2))`
ctx.ReplaceAll([&](ast::TypeConstructorExpression* ctor) -> ast::Expression* {
if (auto* array = As<sem::Array>(sem.Get(ctor)->Type()->UnwrapRef())) {
if (auto w = wrapper(array)) {
// Wrap the array type constructor with another constructor for
// the wrapper
auto* wrapped_array_ty = ctx.Clone(ctor->type());
auto* array_ty = w.array_type(ctx);
auto* arr_ctor =
ctx.dst->Construct(array_ty, ctx.Clone(ctor->values()));
return ctx.dst->Construct(wrapped_array_ty, arr_ctor);
}
}
return nullptr;
});
ctx.Clone();
return Output(Program(std::move(out)));
}
WrapArraysInStructs::WrappedArrayInfo WrapArraysInStructs::WrapArray(
CloneContext& ctx,
std::unordered_map<const sem::Array*, WrappedArrayInfo>& wrapped_arrays,
const sem::Array* array) const {
if (array->IsRuntimeSized()) {
return {}; // We don't want to wrap runtime sized arrays
}
return utils::GetOrCreate(wrapped_arrays, array, [&] {
WrappedArrayInfo info;
// Generate a unique name for the array wrapper
info.wrapper_name = ctx.dst->Symbols().New("tint_array_wrapper");
// Examine the element type. Is it also an array?
std::function<ast::Type*(CloneContext&)> el_type;
if (auto* el_array = array->ElemType()->As<sem::Array>()) {
// Array of array - call WrapArray() on the element type
if (auto el = WrapArray(ctx, wrapped_arrays, el_array)) {
el_type = [=](CloneContext& c) {
return c.dst->create<ast::TypeName>(el.wrapper_name);
};
}
}
// If the element wasn't an array, just create the typical AST type for it
if (!el_type) {
el_type = [=](CloneContext& c) {
return CreateASTTypeFor(&c, array->ElemType());
};
}
// Construct the single structure field type
info.array_type = [=](CloneContext& c) {
ast::DecorationList decos;
if (!array->IsStrideImplicit()) {
decos.emplace_back(
c.dst->create<ast::StrideDecoration>(array->Stride()));
}
return c.dst->create<ast::Array>(el_type(c), array->Count(),
std::move(decos));
};
// 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(info.wrapper_name,
{ctx.dst->Member("arr", info.array_type(ctx))});
return info;
});
}
} // namespace transform
} // namespace tint

84
src/transform/wrap_arrays_in_structs.h Normal file
View File

@ -0,0 +1,84 @@
// 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_TRANSFORM_WRAP_ARRAYS_IN_STRUCTS_H_
#define SRC_TRANSFORM_WRAP_ARRAYS_IN_STRUCTS_H_
#include <string>
#include <unordered_map>
#include "src/transform/transform.h"
// Forward declarations
namespace tint {
namespace ast {
class Type;
} // namespace ast
} // namespace tint
namespace tint {
namespace transform {
/// WrapArraysInStructs is a transform that replaces all array types with a
/// structure holding a single field of that array type.
/// Array index expressions and constructors are also adjusted to deal with this
/// wrapping.
/// This transform helps with backends that cannot directly return arrays or use
/// them as parameters.
class WrapArraysInStructs : public Transform {
public:
/// Constructor
WrapArraysInStructs();
/// Destructor
~WrapArraysInStructs() override;
/// Runs the transform on `program`, returning the transformation result.
/// @param program the source program to transform
/// @param data optional extra transform-specific input data
/// @returns the transformation result
Output Run(const Program* program, const DataMap& data = {}) override;
private:
struct WrappedArrayInfo {
WrappedArrayInfo();
WrappedArrayInfo(const WrappedArrayInfo&);
~WrappedArrayInfo();
Symbol wrapper_name;
std::function<ast::Type*(CloneContext&)> array_type;
operator bool() { return wrapper_name.IsValid(); }
};
/// WrapArray wraps the fixed-size array type in a new structure (if it hasn't
/// already been wrapped). WrapArray will recursively wrap arrays-of-arrays.
/// The new structure will be added to module-scope type declarations of
/// `ctx.dst`.
/// @param ctx the CloneContext
/// @param wrapped_arrays a map of src array type to the wrapped structure
/// name
/// @param array the array type
/// @return the name of the structure that wraps the array, or an invalid
/// Symbol if this array should not be wrapped
WrappedArrayInfo WrapArray(
CloneContext& ctx,
std::unordered_map<const sem::Array*, WrappedArrayInfo>& wrapped_arrays,
const sem::Array* array) const;
};
} // namespace transform
} // namespace tint
#endif // SRC_TRANSFORM_WRAP_ARRAYS_IN_STRUCTS_H_

317
src/transform/wrap_arrays_in_structs_test.cc Normal file
View File

@ -0,0 +1,317 @@
// 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/transform/wrap_arrays_in_structs.h"
#include <memory>
#include <utility>
#include "src/transform/test_helper.h"
namespace tint {
namespace transform {
namespace {
using WrapArraysInStructsTest = TransformTest;
TEST_F(WrapArraysInStructsTest, EmptyModule) {
auto* src = "";
auto* expect = "";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArrayAsGlobal) {
auto* src = R"(
var<private> arr : array<i32, 4>;
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
var<private> arr : tint_array_wrapper;
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArrayAsFunctionVar) {
auto* src = R"(
fn f() {
var arr : array<i32, 4>;
let x = arr[3];
}
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
fn f() {
var arr : tint_array_wrapper;
let x = arr.arr[3];
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArrayAsParam) {
auto* src = R"(
fn f(a : array<i32, 4>) -> i32 {
return a[2];
}
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
fn f(a : tint_array_wrapper) -> i32 {
return a.arr[2];
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArrayAsReturn) {
auto* src = R"(
fn f() -> array<i32, 4> {
return array<i32, 4>(1, 2, 3, 4);
}
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
fn f() -> tint_array_wrapper {
return tint_array_wrapper(array<i32, 4>(1, 2, 3, 4));
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArrayAlias) {
auto* src = R"(
type Inner = array<i32, 2>;
type Array = array<Inner, 2>;
fn f() {
var arr : Array;
arr = Array();
arr = Array(Inner(1, 2), Inner(3, 4));
let vals : Array = Array(Inner(1, 2), Inner(3, 4));
arr = vals;
let x = arr[3];
}
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 2>;
};
type Inner = tint_array_wrapper;
struct tint_array_wrapper_1 {
arr : array<tint_array_wrapper, 2>;
};
type Array = tint_array_wrapper_1;
fn f() {
var arr : tint_array_wrapper_1;
arr = tint_array_wrapper_1(array<tint_array_wrapper, 2>());
arr = tint_array_wrapper_1(array<tint_array_wrapper, 2>(tint_array_wrapper(array<i32, 2>(1, 2)), tint_array_wrapper(array<i32, 2>(3, 4))));
let vals : tint_array_wrapper_1 = tint_array_wrapper_1(array<tint_array_wrapper, 2>(tint_array_wrapper(array<i32, 2>(1, 2)), tint_array_wrapper(array<i32, 2>(3, 4))));
arr = vals;
let x = arr.arr[3];
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArraysInStruct) {
auto* src = R"(
struct S {
a : array<i32, 4>;
b : array<i32, 8>;
c : array<i32, 4>;
};
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
struct tint_array_wrapper_1 {
arr : array<i32, 8>;
};
struct S {
a : tint_array_wrapper;
b : tint_array_wrapper_1;
c : tint_array_wrapper;
};
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, ArraysOfArraysInStruct) {
auto* src = R"(
struct S {
a : array<i32, 4>;
b : array<array<i32, 4>, 4>;
c : array<array<array<i32, 4>, 4>, 4>;
};
)";
auto* expect = R"(
struct tint_array_wrapper {
arr : array<i32, 4>;
};
struct tint_array_wrapper_1 {
arr : array<tint_array_wrapper, 4>;
};
struct tint_array_wrapper_2 {
arr : array<tint_array_wrapper_1, 4>;
};
struct S {
a : tint_array_wrapper;
b : tint_array_wrapper_1;
c : tint_array_wrapper_2;
};
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, AccessArraysOfArraysInStruct) {
auto* src = R"(
struct S {
a : array<i32, 4>;
b : array<array<i32, 4>, 4>;
c : array<array<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_array_wrapper {
arr : array<i32, 4>;
};
struct tint_array_wrapper_1 {
arr : array<tint_array_wrapper, 4>;
};
struct tint_array_wrapper_2 {
arr : array<tint_array_wrapper_1, 4>;
};
struct S {
a : tint_array_wrapper;
b : tint_array_wrapper_1;
c : tint_array_wrapper_2;
};
fn f(s : S) -> i32 {
return ((s.a.arr[2] + s.b.arr[1].arr[2]) + s.c.arr[3].arr[1].arr[2]);
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(WrapArraysInStructsTest, DeclarationOrder) {
auto* src = R"(
type T0 = i32;
type T1 = array<i32, 1>;
type T2 = i32;
fn f1(a : array<i32, 2>) {
}
type T3 = i32;
fn f2() {
var v : array<i32, 3>;
}
)";
auto* expect = R"(
type T0 = i32;
struct tint_array_wrapper {
arr : array<i32, 1>;
};
type T1 = tint_array_wrapper;
type T2 = i32;
struct tint_array_wrapper_1 {
arr : array<i32, 2>;
};
fn f1(a : tint_array_wrapper_1) {
}
type T3 = i32;
struct tint_array_wrapper_2 {
arr : array<i32, 3>;
};
fn f2() {
var v : tint_array_wrapper_2;
}
)";
auto got = Run<WrapArraysInStructs>(src);
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace transform
} // namespace tint

72
src/writer/msl/generator_impl.cc
View File

@ -95,17 +95,7 @@ bool GeneratorImpl::Generate() {
}
}
// Generate wrappers for array types.
for (auto* ty : program_->Types()) {
auto* arr = ty->As<sem::Array>();
if (arr && !arr->IsRuntimeSized()) {
if (!EmitArrayWrapper(arr)) {
return false;
}
}
}
if (!program_->AST().TypeDecls().empty() || array_wrappers_.size()) {
if (!program_->AST().TypeDecls().empty()) {
out_ << std::endl;
}
@ -191,13 +181,6 @@ bool GeneratorImpl::EmitArrayAccessor(ast::ArrayAccessorExpression* expr) {
out_ << ")";
}
// If the array has been wrapped in a struct, emit a member accessor to get to
// the actual array.
auto* sem_ty = program_->Sem().Get(expr->array())->Type()->UnwrapRef();
if (array_wrappers_.count(sem_ty)) {
out_ << ".array";
}
out_ << "[";
if (!EmitExpression(expr->idx_expr())) {
@ -2010,16 +1993,6 @@ bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
if (auto* ary = type->As<sem::Array>()) {
if (array_wrappers_.count(type)) {
// If the array has been wrapped in a struct, emit the struct name instead
// of the array type.
out_ << array_wrappers_[type];
if (!name.empty()) {
out_ << " " << name;
}
return true;
}
const sem::Type* base_type = ary;
std::vector<uint32_t> sizes;
while (auto* arr = base_type->As<sem::Array>()) {
@ -2069,8 +2042,7 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
default:
TINT_ICE(diagnostics_) << "unhandled storage class for pointer";
}
if (ptr->StoreType()->Is<sem::Array>() &&
!array_wrappers_.count(ptr->StoreType())) {
if (ptr->StoreType()->Is<sem::Array>()) {
std::string inner = "(*" + name + ")";
if (!EmitType(ptr->StoreType(), inner)) {
return false;
@ -2185,16 +2157,6 @@ bool GeneratorImpl::EmitPackedType(const sem::Type* type,
}
bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
// Generate wrappers for any struct members that have an array type.
for (auto* mem : str->Members()) {
auto* arr = mem->Type()->As<sem::Array>();
if (arr && !arr->IsRuntimeSized()) {
if (!EmitArrayWrapper(arr)) {
return false;
}
}
}
// TODO(dsinclair): Block decoration?
// if (str->impl()->decoration() != ast::Decoration::kNone) {
// }
@ -2517,36 +2479,6 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
return {};
}
bool GeneratorImpl::EmitArrayWrapper(const sem::Array* arr) {
// We may have created a wrapper for this array while emitting structs.
if (array_wrappers_.count(arr)) {
return true;
}
// Find a unique name for the new structure.
uint32_t i = static_cast<uint32_t>(array_wrappers_.size());
std::string struct_name;
do {
struct_name = "tint_array_wrapper_" + std::to_string(i++);
if (i == 0) {
TINT_ICE(diagnostics_) << "unable to allocate name for array wrapper";
}
} while (program_->Symbols().Get(struct_name).IsValid());
// Generate a struct with a single member with the same type as the array.
out_ << "struct " << struct_name << " {" << std::endl;
out_ << " ";
if (!EmitType(arr->ElemType(), "")) {
return false;
}
out_ << " array[" << arr->Count() << "];" << std::endl;
out_ << "};" << std::endl;
array_wrappers_.emplace(arr, struct_name);
return true;
}
} // namespace msl
} // namespace writer
} // namespace tint

7
src/writer/msl/generator_impl.h
View File

@ -42,7 +42,6 @@ namespace tint {
// Forward declarations
namespace sem {
class Array;
class Call;
class Intrinsic;
} // namespace sem
@ -302,18 +301,12 @@ class GeneratorImpl : public TextGenerator {
/// type.
SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
/// Generate a struct wrapper for an array type.
bool EmitArrayWrapper(const sem::Array* arr);
ScopeStack<const sem::Variable*> global_variables_;
Symbol current_ep_sym_;
bool generating_entry_point_ = false;
const Program* program_ = nullptr;
uint32_t loop_emission_counter_ = 0;
// Map from an array type to the name of a struct which wraps it.
std::unordered_map<const sem::Type*, std::string> array_wrappers_;
std::unordered_map<Symbol, EntryPointData> ep_sym_to_in_data_;
std::unordered_map<Symbol, EntryPointData> ep_sym_to_out_data_;

41
src/writer/msl/generator_impl_function_test.cc
View File

@ -765,12 +765,11 @@ TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) {
params.push_back(Param("a", ty.array<f32, 5>()));
Func("my_func", params, ty.void_(),
ast::StatementList{
{
Return(),
},
{});
});
GeneratorImpl& gen = Build();
GeneratorImpl& gen = SanitizeAndBuild();
gen.increment_indent();
@ -778,17 +777,43 @@ TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayParams) {
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
float array[5];
};
struct tint_array_wrapper {
float arr[5];
};
void my_func(tint_array_wrapper_0 const a) {
void my_func(tint_array_wrapper a) {
return;
}
)");
}
TEST_F(MslGeneratorImplTest, Emit_Function_WithArrayReturn) {
Func("my_func", {}, ty.array<f32, 5>(),
{
Return(Construct(ty.array<f32, 5>())),
});
GeneratorImpl& gen = SanitizeAndBuild();
gen.increment_indent();
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper {
float arr[5];
};
tint_array_wrapper my_func() {
tint_array_wrapper const tint_symbol = {.arr={}};
return tint_symbol;
}
)");
}
// https://crbug.com/tint/297
TEST_F(MslGeneratorImplTest,
Emit_Function_Multiple_EntryPoint_With_Same_ModuleVar) {

33
src/writer/msl/generator_impl_type_test.cc
View File

@ -449,27 +449,20 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayDefaultStride) {
// ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, NAME, SUFFIX)
// for each field of the structure s.
#define ALL_FIELDS() \
FIELD(0x0000, int, a, /*NO SUFFIX*/) \
FIELD(0x0004, tint_array_wrapper_0, b, /*NO SUFFIX*/) \
FIELD(0x0020, float, c, /*NO SUFFIX*/) \
FIELD(0x0024, int8_t, tint_pad_0, [476]) \
FIELD(0x0200, tint_array_wrapper_1, d, /*NO SUFFIX*/) \
FIELD(0x1200, float, e, /*NO SUFFIX*/) \
FIELD(0x1204, float, f, [1]) \
#define ALL_FIELDS() \
FIELD(0x0000, int, a, /*NO SUFFIX*/) \
FIELD(0x0004, float, b, [7]) \
FIELD(0x0020, float, c, /*NO SUFFIX*/) \
FIELD(0x0024, int8_t, tint_pad_0, [476]) \
FIELD(0x0200, inner, d, [4]) \
FIELD(0x1200, float, e, /*NO SUFFIX*/) \
FIELD(0x1204, float, f, [1]) \
FIELD(0x1208, int8_t, tint_pad_1, [504])
// Check that the generated string is as expected.
#define FIELD(ADDR, TYPE, NAME, SUFFIX) \
" /* " #ADDR " */ " #TYPE " " #NAME #SUFFIX ";\n"
auto* expect =
"struct tint_array_wrapper_0 {\n"
" float array[7];\n"
"};\n"
"struct tint_array_wrapper_1 {\n"
" inner array[4];\n"
"};\n"
"struct S {\n" ALL_FIELDS() "};\n";
auto* expect = "struct S {\n" ALL_FIELDS() "};\n";
#undef FIELD
EXPECT_EQ(gen.result(), expect);
@ -489,12 +482,12 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayDefaultStride) {
CHECK_TYPE_SIZE_AND_ALIGN(inner, 1024, 512);
// array_x: size(28), align(4)
using tint_array_wrapper_0 = std::array<float, 7>;
CHECK_TYPE_SIZE_AND_ALIGN(tint_array_wrapper_0, 28, 4);
using array_x = std::array<float, 7>;
CHECK_TYPE_SIZE_AND_ALIGN(array_x, 28, 4);
// array_y: size(4096), align(512)
using tint_array_wrapper_1 = std::array<inner, 4>;
CHECK_TYPE_SIZE_AND_ALIGN(tint_array_wrapper_1, 4096, 512);
using array_y = std::array<inner, 4>;
CHECK_TYPE_SIZE_AND_ALIGN(array_y, 4096, 512);
// array_z: size(4), align(4)
using array_z = std::array<float, 1>;

1
test/BUILD.gn
View File

@ -288,6 +288,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
"../src/transform/transform_test.cc",
"../src/transform/var_for_dynamic_index_test.cc",
"../src/transform/vertex_pulling_test.cc",
"../src/transform/wrap_arrays_in_structs_test.cc",
"../src/utils/enum_set_test.cc",
"../src/utils/get_or_create_test.cc",
"../src/utils/hash_test.cc",

20
test/array/assign_to_subexpr.wgsl.expected.msl
View File

@ -1,28 +1,28 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
int array[4];
struct tint_array_wrapper {
int arr[4];
};
struct S {
tint_array_wrapper_0 arr;
tint_array_wrapper arr;
};
struct tint_array_wrapper_1 {
tint_array_wrapper_0 array[2];
tint_array_wrapper arr[2];
};
void foo() {
tint_array_wrapper_0 const src = {};
tint_array_wrapper_0 dst = {0};
tint_array_wrapper const src = {.arr={}};
tint_array_wrapper dst = {};
S dst_struct = {};
tint_array_wrapper_1 dst_array = {{0}};
thread tint_array_wrapper_0* const dst_ptr = &(dst);
tint_array_wrapper_1 dst_array = {};
thread tint_array_wrapper* const dst_ptr = &(dst);
thread S* const dst_struct_ptr = &(dst_struct);
thread tint_array_wrapper_1* const dst_array_ptr = &(dst_array);
dst_struct.arr = src;
dst_array.array[1] = src;
dst_array.arr[1] = src;
*(dst_ptr) = src;
(*(dst_struct_ptr)).arr = src;
(*(dst_array_ptr)).array[0] = src;
(*(dst_array_ptr)).arr[0] = src;
}

26
test/array/function_parameter.wgsl.expected.msl
View File

@ -1,32 +1,32 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
struct tint_array_wrapper_1 {
tint_array_wrapper_0 array[3];
tint_array_wrapper arr[3];
};
struct tint_array_wrapper_2 {
tint_array_wrapper_1 array[2];
tint_array_wrapper_1 arr[2];
};
float f1(tint_array_wrapper_0 const a) {
return a.array[3];
float f1(tint_array_wrapper a) {
return a.arr[3];
}
float f2(tint_array_wrapper_1 const a) {
return a.array[2].array[3];
float f2(tint_array_wrapper_1 a) {
return a.arr[2].arr[3];
}
float f3(tint_array_wrapper_2 const a) {
return a.array[1].array[2].array[3];
float f3(tint_array_wrapper_2 a) {
return a.arr[1].arr[2].arr[3];
}
kernel void tint_symbol() {
tint_array_wrapper_0 const a1 = {};
tint_array_wrapper_1 const a2 = {};
tint_array_wrapper_2 const a3 = {};
tint_array_wrapper const a1 = {.arr={}};
tint_array_wrapper_1 const a2 = {.arr={}};
tint_array_wrapper_2 const a3 = {.arr={}};
float const v1 = f1(a1);
float const v2 = f2(a2);
float const v3 = f3(a3);

18
test/array/function_return_type.wgsl.expected.msl
View File

@ -1,33 +1,33 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
struct tint_array_wrapper_1 {
tint_array_wrapper_0 array[3];
tint_array_wrapper arr[3];
};
struct tint_array_wrapper_2 {
tint_array_wrapper_1 array[2];
tint_array_wrapper_1 arr[2];
};
tint_array_wrapper_0 f1() {
tint_array_wrapper_0 const tint_symbol_1 = {};
tint_array_wrapper f1() {
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
}
tint_array_wrapper_1 f2() {
tint_array_wrapper_1 const tint_symbol_2 = {f1(), f1(), f1()};
tint_array_wrapper_1 const tint_symbol_2 = {.arr={f1(), f1(), f1()}};
return tint_symbol_2;
}
tint_array_wrapper_2 f3() {
tint_array_wrapper_2 const tint_symbol_3 = {f2(), f2()};
tint_array_wrapper_2 const tint_symbol_3 = {.arr={f2(), f2()}};
return tint_symbol_3;
}
kernel void tint_symbol() {
tint_array_wrapper_0 const a1 = f1();
tint_array_wrapper const a1 = f1();
tint_array_wrapper_1 const a2 = f2();
tint_array_wrapper_2 const a3 = f3();
return;

78
test/array/type_constructor.wgsl.expected.msl
View File

@ -1,53 +1,53 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
int array[4];
};
struct tint_array_wrapper_1 {
tint_array_wrapper_0 array[3];
struct tint_array_wrapper {
int arr[4];
};
struct tint_array_wrapper_2 {
tint_array_wrapper_1 array[2];
tint_array_wrapper arr[3];
};
struct tint_array_wrapper_1 {
tint_array_wrapper_2 arr[2];
};
struct tint_array_wrapper_3 {
tint_array_wrapper_0 array[2];
tint_array_wrapper arr[2];
};
kernel void tint_symbol() {
int const x = 42;
tint_array_wrapper_0 const empty = {};
tint_array_wrapper_0 const nonempty = {1, 2, 3, 4};
tint_array_wrapper_0 const nonempty_with_expr = {1, x, (x + 1), nonempty.array[3]};
tint_array_wrapper_2 const nested_empty = {};
tint_array_wrapper_0 const tint_symbol_1 = {1, 2, 3, 4};
tint_array_wrapper_0 const tint_symbol_2 = {5, 6, 7, 8};
tint_array_wrapper_0 const tint_symbol_3 = {9, 10, 11, 12};
tint_array_wrapper_1 const tint_symbol_4 = {tint_symbol_1, tint_symbol_2, tint_symbol_3};
tint_array_wrapper_0 const tint_symbol_5 = {13, 14, 15, 16};
tint_array_wrapper_0 const tint_symbol_6 = {17, 18, 19, 20};
tint_array_wrapper_0 const tint_symbol_7 = {21, 22, 23, 24};
tint_array_wrapper_1 const tint_symbol_8 = {tint_symbol_5, tint_symbol_6, tint_symbol_7};
tint_array_wrapper_2 const nested_nonempty = {tint_symbol_4, tint_symbol_8};
tint_array_wrapper_0 const tint_symbol_9 = {1, 2, x, (x + 1)};
tint_array_wrapper_0 const tint_symbol_10 = {5, 6, nonempty.array[2], (nonempty.array[3] + 1)};
tint_array_wrapper_1 const tint_symbol_11 = {tint_symbol_9, tint_symbol_10, nonempty};
tint_array_wrapper_2 const nested_nonempty_with_expr = {tint_symbol_11, nested_nonempty.array[1]};
tint_array_wrapper_0 const tint_symbol_12 = {};
int const subexpr_empty = tint_symbol_12.array[1];
tint_array_wrapper_0 const tint_symbol_13 = {1, 2, 3, 4};
int const subexpr_nonempty = tint_symbol_13.array[2];
tint_array_wrapper_0 const tint_symbol_14 = {1, x, (x + 1), nonempty.array[3]};
int const subexpr_nonempty_with_expr = tint_symbol_14.array[2];
tint_array_wrapper_3 const tint_symbol_15 = {};
tint_array_wrapper_0 const subexpr_nested_empty = tint_symbol_15.array[1];
tint_array_wrapper_0 const tint_symbol_16 = {1, 2, 3, 4};
tint_array_wrapper_0 const tint_symbol_17 = {5, 6, 7, 8};
tint_array_wrapper_3 const tint_symbol_18 = {tint_symbol_16, tint_symbol_17};
tint_array_wrapper_0 const subexpr_nested_nonempty = tint_symbol_18.array[1];
tint_array_wrapper_0 const tint_symbol_19 = {1, x, (x + 1), nonempty.array[3]};
tint_array_wrapper_3 const tint_symbol_20 = {tint_symbol_19, nested_nonempty.array[1].array[2]};
tint_array_wrapper_0 const subexpr_nested_nonempty_with_expr = tint_symbol_20.array[1];
tint_array_wrapper const empty = {.arr={}};
tint_array_wrapper const nonempty = {.arr={1, 2, 3, 4}};
tint_array_wrapper const nonempty_with_expr = {.arr={1, x, (x + 1), nonempty.arr[3]}};
tint_array_wrapper_1 const nested_empty = {.arr={}};
tint_array_wrapper const tint_symbol_1 = {.arr={1, 2, 3, 4}};
tint_array_wrapper const tint_symbol_2 = {.arr={5, 6, 7, 8}};
tint_array_wrapper const tint_symbol_3 = {.arr={9, 10, 11, 12}};
tint_array_wrapper_2 const tint_symbol_4 = {.arr={tint_symbol_1, tint_symbol_2, tint_symbol_3}};
tint_array_wrapper const tint_symbol_5 = {.arr={13, 14, 15, 16}};
tint_array_wrapper const tint_symbol_6 = {.arr={17, 18, 19, 20}};
tint_array_wrapper const tint_symbol_7 = {.arr={21, 22, 23, 24}};
tint_array_wrapper_2 const tint_symbol_8 = {.arr={tint_symbol_5, tint_symbol_6, tint_symbol_7}};
tint_array_wrapper_1 const nested_nonempty = {.arr={tint_symbol_4, tint_symbol_8}};
tint_array_wrapper const tint_symbol_9 = {.arr={1, 2, x, (x + 1)}};
tint_array_wrapper const tint_symbol_10 = {.arr={5, 6, nonempty.arr[2], (nonempty.arr[3] + 1)}};
tint_array_wrapper_2 const tint_symbol_11 = {.arr={tint_symbol_9, tint_symbol_10, nonempty}};
tint_array_wrapper_1 const nested_nonempty_with_expr = {.arr={tint_symbol_11, nested_nonempty.arr[1]}};
tint_array_wrapper const tint_symbol_12 = {.arr={}};
int const subexpr_empty = tint_symbol_12.arr[1];
tint_array_wrapper const tint_symbol_13 = {.arr={1, 2, 3, 4}};
int const subexpr_nonempty = tint_symbol_13.arr[2];
tint_array_wrapper const tint_symbol_14 = {.arr={1, x, (x + 1), nonempty.arr[3]}};
int const subexpr_nonempty_with_expr = tint_symbol_14.arr[2];
tint_array_wrapper_3 const tint_symbol_15 = {.arr={}};
tint_array_wrapper const subexpr_nested_empty = tint_symbol_15.arr[1];
tint_array_wrapper const tint_symbol_16 = {.arr={1, 2, 3, 4}};
tint_array_wrapper const tint_symbol_17 = {.arr={5, 6, 7, 8}};
tint_array_wrapper_3 const tint_symbol_18 = {.arr={tint_symbol_16, tint_symbol_17}};
tint_array_wrapper const subexpr_nested_nonempty = tint_symbol_18.arr[1];
tint_array_wrapper const tint_symbol_19 = {.arr={1, x, (x + 1), nonempty.arr[3]}};
tint_array_wrapper_3 const tint_symbol_20 = {.arr={tint_symbol_19, nested_nonempty.arr[1].arr[2]}};
tint_array_wrapper const subexpr_nested_nonempty_with_expr = tint_symbol_20.arr[1];
return;
}

8
test/bug/tint/782.wgsl.expected.msl
View File

@ -1,13 +1,13 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper_0 {
int array[2];
struct tint_array_wrapper {
int arr[2];
};
void foo() {
tint_array_wrapper_0 tint_symbol = {0};
tint_array_wrapper_0 implict = {0};
tint_array_wrapper tint_symbol = {};
tint_array_wrapper implict = {};
implict = tint_symbol;
}

14
test/bug/tint/824.wgsl.expected.msl
View File

@ -9,20 +9,20 @@ struct tint_symbol_2 {
float4 color [[user(locn0)]];
float4 Position [[position]];
};
struct tint_array_wrapper_0 {
float2 array[4];
struct tint_array_wrapper {
float2 arr[4];
};
struct tint_array_wrapper_1 {
float4 array[4];
float4 arr[4];
};
vertex tint_symbol_2 tint_symbol(uint VertexIndex [[vertex_id]], uint InstanceIndex [[instance_id]]) {
tint_array_wrapper_0 const zv = {float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)};
float const z = zv.array[InstanceIndex].x;
tint_array_wrapper const zv = {.arr={float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)}};
float const z = zv.arr[InstanceIndex].x;
Output output = {};
output.Position = float4(0.5f, 0.5f, z, 1.0f);
tint_array_wrapper_1 const colors = {float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)};
output.color = colors.array[InstanceIndex];
tint_array_wrapper_1 const colors = {.arr={float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)}};
output.color = colors.arr[InstanceIndex];
tint_symbol_2 const tint_symbol_3 = {.color=output.color, .Position=output.Position};
return tint_symbol_3;
}

18
test/samples/compute_boids.wgsl.expected.msl
View File

@ -25,11 +25,11 @@ struct SimParams {
/* 0x0014 */ float rule2Scale;
/* 0x0018 */ float rule3Scale;
};
struct tint_array_wrapper_0 {
Particle array[5];
struct tint_array_wrapper {
/* 0x0000 */ Particle arr[5];
};
struct Particles {
/* 0x0000 */ tint_array_wrapper_0 particles;
/* 0x0000 */ tint_array_wrapper particles;
};
vertex tint_symbol_2 vert_main(tint_symbol_1 tint_symbol [[stage_in]]) {
@ -52,8 +52,8 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
if ((index >= 5u)) {
return;
}
float2 vPos = particlesA.particles.array[index].pos;
float2 vVel = particlesA.particles.array[index].vel;
float2 vPos = particlesA.particles.arr[index].pos;
float2 vVel = particlesA.particles.arr[index].vel;
float2 cMass = float2(0.0f, 0.0f);
float2 cVel = float2(0.0f, 0.0f);
float2 colVel = float2(0.0f, 0.0f);
@ -77,8 +77,8 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
if ((i == index)) {
continue;
}
pos = particlesA.particles.array[i].pos.xy;
vel = particlesA.particles.array[i].vel.xy;
pos = particlesA.particles.arr[i].pos.xy;
vel = particlesA.particles.arr[i].vel.xy;
if (( distance(pos, vPos) < params.rule1Distance)) {
cMass = (cMass + pos);
cMassCount = (cMassCount + 1);
@ -114,8 +114,8 @@ kernel void comp_main(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], c
if ((vPos.y > 1.0f)) {
vPos.y = -1.0f;
}
particlesB.particles.array[index].pos = vPos;
particlesB.particles.array[index].vel = vVel;
particlesB.particles.arr[index].pos = vPos;
particlesB.particles.arr[index].vel = vVel;
return;
}

10
test/samples/triangle.wgsl.expected.msl
View File

@ -1,19 +1,19 @@
#include <metal_stdlib>
using namespace metal;
struct tint_array_wrapper {
float2 arr[3];
};
struct tint_symbol_1 {
float4 value [[position]];
};
struct tint_symbol_2 {
float4 value [[color(0)]];
};
struct tint_array_wrapper_0 {
float2 array[3];
};
constant tint_array_wrapper_0 pos = {float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)};
constant tint_array_wrapper pos = {.arr={float2(0.0f, 0.5f), float2(-0.5f, -0.5f), float2(0.5f, -0.5f)}};
vertex tint_symbol_1 vtx_main(uint VertexIndex [[vertex_id]]) {
tint_symbol_1 const tint_symbol_3 = {.value=float4(pos.array[VertexIndex], 0.0f, 1.0f)};
tint_symbol_1 const tint_symbol_3 = {.value=float4(pos.arr[VertexIndex], 0.0f, 1.0f)};
return tint_symbol_3;
}

20
test/struct/type_constructor.wgsl.expected.msl
View File

@ -16,14 +16,14 @@ struct S3 {
S1 h;
S2 i;
};
struct tint_array_wrapper_0 {
int array[2];
struct tint_array_wrapper {
int arr[2];
};
struct T {
tint_array_wrapper_0 a;
tint_array_wrapper a;
};
struct tint_array_wrapper_1 {
T array[2];
T arr[2];
};
kernel void tint_symbol() {
@ -53,15 +53,15 @@ kernel void tint_symbol() {
S1 const tint_symbol_12 = {.a=2, .b=x, .c=(x + 1), .d=nested_nonempty.i.f.d};
S2 const tint_symbol_13 = {.e=1, .f=tint_symbol_12};
S1 const subexpr_nested_nonempty_with_expr = tint_symbol_13.f;
tint_array_wrapper_1 const aosoa_empty = {};
tint_array_wrapper_0 const tint_symbol_14 = {1, 2};
tint_array_wrapper_1 const aosoa_empty = {.arr={}};
tint_array_wrapper const tint_symbol_14 = {.arr={1, 2}};
T const tint_symbol_15 = {.a=tint_symbol_14};
tint_array_wrapper_0 const tint_symbol_16 = {3, 4};
tint_array_wrapper const tint_symbol_16 = {.arr={3, 4}};
T const tint_symbol_17 = {.a=tint_symbol_16};
tint_array_wrapper_1 const aosoa_nonempty = {tint_symbol_15, tint_symbol_17};
tint_array_wrapper_0 const tint_symbol_18 = {1, (aosoa_nonempty.array[0].a.array[0] + 1)};
tint_array_wrapper_1 const aosoa_nonempty = {.arr={tint_symbol_15, tint_symbol_17}};
tint_array_wrapper const tint_symbol_18 = {.arr={1, (aosoa_nonempty.arr[0].a.arr[0] + 1)}};
T const tint_symbol_19 = {.a=tint_symbol_18};
tint_array_wrapper_1 const aosoa_nonempty_with_expr = {tint_symbol_19, aosoa_nonempty.array[1]};
tint_array_wrapper_1 const aosoa_nonempty_with_expr = {.arr={tint_symbol_19, aosoa_nonempty.arr[1]}};
return;
}

10
test/types/function_scope_declarations.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
kernel void tint_symbol() {
@ -24,13 +24,13 @@ kernel void tint_symbol() {
float4 const v4f32_let = float4();
float2x3 m2x3_var = float2x3();
float3x4 const m3x4_let = float3x4();
tint_array_wrapper_0 arr_var = {};
tint_array_wrapper_0 const arr_let = {};
tint_array_wrapper arr_var = {.arr={}};
tint_array_wrapper const arr_let = {.arr={}};
S struct_var = {};
S const struct_let = {};
thread float* const ptr_f32 = &(f32_var);
thread float4* const ptr_vec = &(v4f32_var);
thread tint_array_wrapper_0* const ptr_arr = &(arr_var);
thread tint_array_wrapper* const ptr_arr = &(arr_var);
return;
}

6
test/types/module_scope_let.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
constant bool bool_let = bool();
@ -15,7 +15,7 @@ constant int2 v2i32_let = int2();
constant uint3 v3u32_let = uint3();
constant float4 v4f32_let = float4();
constant float3x4 m3x4_let = float3x4();
constant tint_array_wrapper_0 arr_let = {};
constant tint_array_wrapper arr_let = {.arr={}};
constant S struct_let = {};
kernel void tint_symbol() {
return;

8
test/types/module_scope_var.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
kernel void tint_symbol() {
@ -16,7 +16,7 @@ kernel void tint_symbol() {
thread uint3 tint_symbol_8 = 0u;
thread float4 tint_symbol_9 = 0.0f;
thread float2x3 tint_symbol_10 = float2x3(0.0f);
thread tint_array_wrapper_0 tint_symbol_11 = {0.0f};
thread tint_array_wrapper tint_symbol_11 = {};
thread S tint_symbol_12 = {};
tint_symbol_3 = bool();
tint_symbol_4 = int();
@ -26,7 +26,7 @@ kernel void tint_symbol() {
tint_symbol_8 = uint3();
tint_symbol_9 = float4();
tint_symbol_10 = float2x3();
tint_array_wrapper_0 const tint_symbol_1 = {};
tint_array_wrapper const tint_symbol_1 = {.arr={}};
tint_symbol_11 = tint_symbol_1;
S const tint_symbol_2 = {};
tint_symbol_12 = tint_symbol_2;

8
test/types/module_scope_var_initializers.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
kernel void tint_symbol() {
@ -16,7 +16,7 @@ kernel void tint_symbol() {
thread uint3 tint_symbol_8 = uint3();
thread float4 tint_symbol_9 = float4();
thread float2x3 tint_symbol_10 = float2x3();
thread tint_array_wrapper_0 tint_symbol_11 = {};
thread tint_array_wrapper tint_symbol_11 = {.arr={}};
thread S tint_symbol_12 = {};
tint_symbol_3 = bool();
tint_symbol_4 = int();
@ -26,7 +26,7 @@ kernel void tint_symbol() {
tint_symbol_8 = uint3();
tint_symbol_9 = float4();
tint_symbol_10 = float2x3();
tint_array_wrapper_0 const tint_symbol_1 = {};
tint_array_wrapper const tint_symbol_1 = {.arr={}};
tint_symbol_11 = tint_symbol_1;
S const tint_symbol_2 = {};
tint_symbol_12 = tint_symbol_2;

6
test/types/parameters.wgsl.expected.msl
View File

@ -3,11 +3,11 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
void foo(bool param_bool, int param_i32, uint param_u32, float param_f32, int2 param_v2i32, uint3 param_v3u32, float4 param_v4f32, float2x3 param_m2x3, tint_array_wrapper_0 const param_arr, S param_struct, thread float* const param_ptr_f32, thread float4* const param_ptr_vec, thread tint_array_wrapper_0* const param_ptr_arr) {
void foo(bool param_bool, int param_i32, uint param_u32, float param_f32, int2 param_v2i32, uint3 param_v3u32, float4 param_v4f32, float2x3 param_m2x3, tint_array_wrapper param_arr, S param_struct, thread float* const param_ptr_f32, thread float4* const param_ptr_vec, thread tint_array_wrapper* const param_ptr_arr) {
}
kernel void tint_symbol() {

8
test/types/return_types.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
bool ret_bool() {
@ -39,8 +39,8 @@ float2x3 ret_m2x3() {
return float2x3();
}
tint_array_wrapper_0 ret_arr() {
tint_array_wrapper_0 const tint_symbol_1 = {};
tint_array_wrapper ret_arr() {
tint_array_wrapper const tint_symbol_1 = {.arr={}};
return tint_symbol_1;
}

6
test/types/struct_members.wgsl.expected.msl
View File

@ -3,8 +3,8 @@
using namespace metal;
struct S_inner {
};
struct tint_array_wrapper_0 {
float array[4];
struct tint_array_wrapper {
float arr[4];
};
struct S {
bool member_bool;
@ -15,7 +15,7 @@ struct S {
uint3 member_v3u32;
float4 member_v4f32;
float2x3 member_m2x3;
tint_array_wrapper_0 member_arr;
tint_array_wrapper member_arr;
S_inner member_struct;
};

14
test/var/inferred/function-let.wgsl.expected.msl
View File

@ -4,12 +4,12 @@ using namespace metal;
struct MyStruct {
float f1;
};
struct tint_array_wrapper {
float arr[10];
};
struct tint_symbol_1 {
float4 value [[color(0)]];
};
struct tint_array_wrapper_0 {
float array[10];
};
int ret_i32() {
return 1;
@ -28,8 +28,8 @@ MyStruct ret_MyStruct() {
return tint_symbol_2;
}
tint_array_wrapper_0 ret_MyArray() {
tint_array_wrapper_0 const tint_symbol_3 = {};
tint_array_wrapper ret_MyArray() {
tint_array_wrapper const tint_symbol_3 = {.arr={}};
return tint_symbol_3;
}
@ -42,13 +42,13 @@ void let_decls() {
float3 const v6 = float3(1.0f, 1.0f, 1.0f);
float3x3 const v7 = float3x3(v6, v6, v6);
MyStruct const v8 = {.f1=1.0f};
tint_array_wrapper_0 const v9 = {};
tint_array_wrapper const v9 = {.arr={}};
int const v10 = ret_i32();
uint const v11 = ret_u32();
float const v12 = ret_f32();
MyStruct const v13 = ret_MyStruct();
MyStruct const v14 = ret_MyStruct();
tint_array_wrapper_0 const v15 = ret_MyArray();
tint_array_wrapper const v15 = ret_MyArray();
}
fragment tint_symbol_1 tint_symbol() {

14
test/var/inferred/function-var.wgsl.expected.msl
View File

@ -4,12 +4,12 @@ using namespace metal;
struct MyStruct {
float f1;
};
struct tint_array_wrapper {
float arr[10];
};
struct tint_symbol_1 {
float4 value [[color(0)]];
};
struct tint_array_wrapper_0 {
float array[10];
};
int ret_i32() {
return 1;
@ -28,8 +28,8 @@ MyStruct ret_MyStruct() {
return tint_symbol_2;
}
tint_array_wrapper_0 ret_MyArray() {
tint_array_wrapper_0 const tint_symbol_3 = {};
tint_array_wrapper ret_MyArray() {
tint_array_wrapper const tint_symbol_3 = {.arr={}};
return tint_symbol_3;
}
@ -42,13 +42,13 @@ void var_decls() {
float3 v6 = float3(1.0f, 1.0f, 1.0f);
float3x3 v7 = float3x3(v6, v6, v6);
MyStruct v8 = {.f1=1.0f};
tint_array_wrapper_0 v9 = {};
tint_array_wrapper v9 = {.arr={}};
int v10 = ret_i32();
uint v11 = ret_u32();
float v12 = ret_f32();
MyStruct v13 = ret_MyStruct();
MyStruct v14 = ret_MyStruct();
tint_array_wrapper_0 v15 = ret_MyArray();
tint_array_wrapper v15 = ret_MyArray();
}
fragment tint_symbol_1 tint_symbol() {

8
test/var/inferred/global-let.wgsl.expected.msl
View File

@ -4,12 +4,12 @@ using namespace metal;
struct MyStruct {
float f1;
};
struct tint_array_wrapper {
float arr[10];
};
struct tint_symbol_1 {
float4 value [[color(0)]];
};
struct tint_array_wrapper_0 {
float array[10];
};
constant int v1 = 1;
constant uint v2 = 1u;
@ -19,7 +19,7 @@ constant uint3 v5 = uint3(1u, 1u, 1u);
constant float3 v6 = float3(1.0f, 1.0f, 1.0f);
constant float3x3 v7 = float3x3(float3(1.0f, 1.0f, 1.0f), float3(1.0f, 1.0f, 1.0f), float3(1.0f, 1.0f, 1.0f));
constant MyStruct v8 = {};
constant tint_array_wrapper_0 v9 = {};
constant tint_array_wrapper v9 = {.arr={}};
fragment tint_symbol_1 tint_symbol() {
tint_symbol_1 const tint_symbol_2 = {.value=float4(0.0f, 0.0f, 0.0f, 0.0f)};
return tint_symbol_2;