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

Deprecate the @stride attribute 

Update validation error for invalid uniform array element alignment.

Update tests to either remove the @stride attribute or use a different
element type.

Bug: tint:1381
Change-Id: I50b52cd78a34d9cd162fa5f2171a5fd35dcf3b79
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/77560
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
James Price 2022-01-20 22:11:07 +00:00
parent e04d0f40de
commit f6e5cc03bf
71 changed files with 575 additions and 512 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
docs/origin-trial-changes.md
View File

@ -13,6 +13,7 @@ The following features have been deprecated and will be removed in M102:
* The `[[block]]` attribute has been deprecated. [tint:1324](https://crbug.com/tint/1324) * The `[[block]]` attribute has been deprecated. [tint:1324](https://crbug.com/tint/1324)
* Attributes now use the `@decoration` syntax instead of the `[[decoration]]` syntax. [tint:1382](https://crbug.com/tint/1382) * Attributes now use the `@decoration` syntax instead of the `[[decoration]]` syntax. [tint:1382](https://crbug.com/tint/1382)
* `elseif` has been replaced with `else if`. [tint:1289](https://crbug.com/tint/1289) * `elseif` has been replaced with `else if`. [tint:1289](https://crbug.com/tint/1289)
* The `[[stride]]` attribute has been deprecated. [tint:1381](https://crbug.com/tint/1381)
### New Features ### New Features

4
src/reader/wgsl/parser_impl.cc
View File

@ -3076,7 +3076,9 @@ Maybe<const ast::Decoration*> ParserImpl::decoration() {
auto val = expect_nonzero_positive_sint(use); auto val = expect_nonzero_positive_sint(use);
if (val.errored) if (val.errored)
return Failure::kErrored; return Failure::kErrored;
deprecated(t.source(),
"the @stride attribute is deprecated; use a larger type if "
"necessary");
return create<ast::StrideDecoration>(t.source(), val.value); return create<ast::StrideDecoration>(t.source(), val.value);
}); });
} }

26
src/reader/wgsl/parser_impl_error_msg_test.cc
View File

@ -985,24 +985,24 @@ var i : array<u32, 3;
} }
TEST_F(ParserImplErrorTest, GlobalDeclVarArrayDecoNotArray) { TEST_F(ParserImplErrorTest, GlobalDeclVarArrayDecoNotArray) {
EXPECT("var i : @stride(1) i32;", EXPECT("var i : @location(1) i32;",
R"(test.wgsl:1:10 error: unexpected decorations R"(test.wgsl:1:10 error: unexpected decorations
var i : @stride(1) i32; var i : @location(1) i32;
^^^^^^ ^^^^^^^^
)"); )");
} }
// TODO(crbug.com/tint/1382): Remove // TODO(crbug.com/tint/1382): Remove
TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayDecoMissingEnd) { TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayDecoMissingEnd) {
EXPECT( EXPECT(
"var i : [[stride(1) array<i32>;", "var i : [[location(1) array<i32>;",
R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
var i : [[stride(1) array<i32>; var i : [[location(1) array<i32>;
^^ ^^
test.wgsl:1:21 error: expected ']]' for decoration list test.wgsl:1:23 error: expected ']]' for decoration list
var i : [[stride(1) array<i32>; var i : [[location(1) array<i32>;
^^^^^ ^^^^^
)"); )");
} }
@ -1025,14 +1025,14 @@ var i : [[stride 1)]] array<i32>;
TEST_F(ParserImplErrorTest, TEST_F(ParserImplErrorTest,
DEPRECATED_GlobalDeclVarArrayDecoStrideMissingRParen) { DEPRECATED_GlobalDeclVarArrayDecoStrideMissingRParen) {
EXPECT( EXPECT(
"var i : [[stride(1]] array<i32>;", "var i : [[location(1]] array<i32>;",
R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
var i : [[stride(1]] array<i32>; var i : [[location(1]] array<i32>;
^^ ^^
test.wgsl:1:19 error: expected ')' for stride decoration test.wgsl:1:21 error: expected ')' for location decoration
var i : [[stride(1]] array<i32>; var i : [[location(1]] array<i32>;
^^ ^^
)"); )");
} }

13
src/reader/wgsl/parser_impl_type_decl_test.cc
View File

@ -534,7 +534,10 @@ TEST_F(ParserImplTest, TypeDecl_Array_Decoration_MissingArray) {
EXPECT_FALSE(t.matched); EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr); ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), "1:2: unexpected decorations"); EXPECT_EQ(
p->error(),
R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
1:2: unexpected decorations)");
} }
TEST_F(ParserImplTest, TypeDecl_Array_Decoration_UnknownDecoration) { TEST_F(ParserImplTest, TypeDecl_Array_Decoration_UnknownDecoration) {
@ -546,6 +549,7 @@ TEST_F(ParserImplTest, TypeDecl_Array_Decoration_UnknownDecoration) {
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:2: expected decoration)"); EXPECT_EQ(p->error(), R"(1:2: expected decoration)");
} }
TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) { TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) {
auto p = parser("@stride 4) array<f32, 5>"); auto p = parser("@stride 4) array<f32, 5>");
auto t = p->type_decl(); auto t = p->type_decl();
@ -563,7 +567,10 @@ TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingRightParen) {
EXPECT_FALSE(t.matched); EXPECT_FALSE(t.matched);
ASSERT_EQ(t.value, nullptr); ASSERT_EQ(t.value, nullptr);
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:11: expected ')' for stride decoration)"); EXPECT_EQ(
p->error(),
R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
1:11: expected ')' for stride decoration)");
} }
TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingValue) { TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingValue) {
@ -610,6 +617,7 @@ TEST_F(ParserImplTest,
EXPECT_EQ( EXPECT_EQ(
p->error(), p->error(),
R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
1:14: expected ']]' for decoration list)"); 1:14: expected ']]' for decoration list)");
} }
@ -638,6 +646,7 @@ TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_MissingRightParen) {
EXPECT_EQ( EXPECT_EQ(
p->error(), p->error(),
R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
1:11: expected ')' for stride decoration)"); 1:11: expected ')' for stride decoration)");
} }

14
src/reader/wgsl/parser_impl_variable_ident_decl_test.cc
View File

@ -77,7 +77,7 @@ TEST_F(ParserImplTest, VariableIdentDecl_InvalidIdent) {
} }
TEST_F(ParserImplTest, VariableIdentDecl_NonAccessDecoFail) { TEST_F(ParserImplTest, VariableIdentDecl_NonAccessDecoFail) {
auto p = parser("my_var : @stride(1) S"); auto p = parser("my_var : @location(1) S");
auto* mem = Member("a", ty.i32(), ast::DecorationList{}); auto* mem = Member("a", ty.i32(), ast::DecorationList{});
ast::StructMemberList members; ast::StructMemberList members;
@ -94,11 +94,11 @@ TEST_F(ParserImplTest, VariableIdentDecl_NonAccessDecoFail) {
} }
TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingRightParen) { TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingRightParen) {
auto p = parser("my_var : @stride(4 S"); auto p = parser("my_var : @location(4 S");
auto decl = p->expect_variable_ident_decl("test"); auto decl = p->expect_variable_ident_decl("test");
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored); ASSERT_TRUE(decl.errored);
ASSERT_EQ(p->error(), "1:20: expected ')' for stride decoration"); ASSERT_EQ(p->error(), "1:22: expected ')' for location decoration");
} }
TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingLeftParen) { TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingLeftParen) {
@ -112,27 +112,27 @@ TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingLeftParen) {
// TODO(crbug.com/tint/1382): Remove // TODO(crbug.com/tint/1382): Remove
TEST_F(ParserImplTest, TEST_F(ParserImplTest,
DEPRECATED_VariableIdentDecl_DecorationMissingRightBlock) { DEPRECATED_VariableIdentDecl_DecorationMissingRightBlock) {
auto p = parser("my_var : [[stride(4) S"); auto p = parser("my_var : [[location(4) S");
auto decl = p->expect_variable_ident_decl("test"); auto decl = p->expect_variable_ident_decl("test");
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored); ASSERT_TRUE(decl.errored);
ASSERT_EQ( ASSERT_EQ(
p->error(), p->error(),
R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
1:22: expected ']]' for decoration list)"); 1:24: expected ']]' for decoration list)");
} }
// TODO(crbug.com/tint/1382): Remove // TODO(crbug.com/tint/1382): Remove
TEST_F(ParserImplTest, TEST_F(ParserImplTest,
DEPRECATED_VariableIdentDecl_DecorationMissingRightParen) { DEPRECATED_VariableIdentDecl_DecorationMissingRightParen) {
auto p = parser("my_var : [[stride(4]] S"); auto p = parser("my_var : [[location(4]] S");
auto decl = p->expect_variable_ident_decl("test"); auto decl = p->expect_variable_ident_decl("test");
ASSERT_TRUE(p->has_error()); ASSERT_TRUE(p->has_error());
ASSERT_TRUE(decl.errored); ASSERT_TRUE(decl.errored);
ASSERT_EQ( ASSERT_EQ(
p->error(), p->error(),
R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
1:20: expected ')' for stride decoration)"); 1:22: expected ')' for location decoration)");
} }
// TODO(crbug.com/tint/1382): Remove // TODO(crbug.com/tint/1382): Remove

26
src/resolver/resolver_validation.cc
View File

@ -321,15 +321,29 @@ bool Resolver::ValidateStorageClassLayout(const sem::Struct* str,
// bytes above, so we only need to validate that stride is a multiple of // bytes above, so we only need to validate that stride is a multiple of
// 16 bytes. // 16 bytes.
if (arr->Stride() % 16 != 0) { if (arr->Stride() % 16 != 0) {
// Since WGSL has no stride attribute, try to provide a useful hint
// for how the shader author can resolve the issue.
std::string hint;
if (arr->ElemType()->is_scalar()) {
hint =
"Consider using a vector or struct as the element type "
"instead.";
} else if (auto* vec = arr->ElemType()->As<sem::Vector>();
vec && vec->type()->Size() == 4) {
hint = "Consider using a vec4 instead.";
} else if (arr->ElemType()->Is<sem::Struct>()) {
hint =
"Consider using the @size attribute on the last struct member.";
} else {
hint =
"Consider wrapping the element type in a struct and using the "
"@size attribute.";
}
AddError( AddError(
"uniform storage requires that array elements be aligned to 16 " "uniform storage requires that array elements be aligned to 16 "
"bytes, but array stride of '" + "bytes, but array element alignment of '" +
member_name_of(m) + "' is currently " + member_name_of(m) + "' is currently " +
std::to_string(arr->Stride()) + std::to_string(arr->Stride()) + ". " + hint,
". Consider setting @stride(" +
std::to_string(
utils::RoundUp(required_align, arr->Stride())) +
") on the array type",
m->Declaration()->type->source); m->Declaration()->type->source);
AddNote("see layout of struct:\n" + str->Layout(builder_->Symbols()), AddNote("see layout of struct:\n" + str->Layout(builder_->Symbols()),
str->Declaration()->source); str->Declaration()->source);

93
src/resolver/storage_class_layout_validation_test.cc
View File

@ -378,8 +378,8 @@ TEST_F(ResolverStorageClassLayoutValidationTest,
// Detect array stride must be a multiple of 16 bytes for uniform buffers // Detect array stride must be a multiple of 16 bytes for uniform buffers
TEST_F(ResolverStorageClassLayoutValidationTest, TEST_F(ResolverStorageClassLayoutValidationTest,
UniformBuffer_InvalidArrayStride) { UniformBuffer_InvalidArrayStride_Scalar) {
// type Inner = @stride(8) array<f32, 10>; // type Inner = array<f32, 10>;
// //
// [[block]] // [[block]]
// struct Outer { // struct Outer {
@ -390,7 +390,7 @@ TEST_F(ResolverStorageClassLayoutValidationTest,
// @group(0) @binding(0) // @group(0) @binding(0)
// var<uniform> a : Outer; // var<uniform> a : Outer;
Alias("Inner", ty.array(ty.f32(), 10, 8)); Alias("Inner", ty.array(ty.f32(), 10));
Structure(Source{{12, 34}}, "Outer", Structure(Source{{12, 34}}, "Outer",
{ {
@ -405,10 +405,93 @@ TEST_F(ResolverStorageClassLayoutValidationTest,
ASSERT_FALSE(r()->Resolve()); ASSERT_FALSE(r()->Resolve());
EXPECT_EQ( EXPECT_EQ(
r()->error(), r()->error(),
R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array stride of 'inner' is currently 8. Consider setting @stride(16) on the array type R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 4. Consider using a vector or struct as the element type instead.
12:34 note: see layout of struct:
/* align(4) size(44) */ struct Outer {
/* offset( 0) align(4) size(40) */ inner : array<f32, 10>;
/* offset(40) align(4) size( 4) */ scalar : i32;
/* */ };
78:90 note: see declaration of variable)");
}
TEST_F(ResolverStorageClassLayoutValidationTest,
UniformBuffer_InvalidArrayStride_Vector) {
// type Inner = array<vec2<f32>, 10>;
//
// [[block]]
// struct Outer {
// inner : Inner;
// scalar : i32;
// };
//
// @group(0) @binding(0)
// var<uniform> a : Outer;
Alias("Inner", ty.array(ty.vec2<f32>(), 10));
Structure(Source{{12, 34}}, "Outer",
{
Member("inner", ty.type_name(Source{{34, 56}}, "Inner")),
Member("scalar", ty.i32()),
},
{StructBlock()});
Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
ast::StorageClass::kUniform, GroupAndBinding(0, 0));
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 8. Consider using a vec4 instead.
12:34 note: see layout of struct:
/* align(8) size(88) */ struct Outer {
/* offset( 0) align(8) size(80) */ inner : array<vec2<f32>, 10>;
/* offset(80) align(4) size( 4) */ scalar : i32;
/* offset(84) align(1) size( 4) */ // -- implicit struct size padding --;
/* */ };
78:90 note: see declaration of variable)");
}
TEST_F(ResolverStorageClassLayoutValidationTest,
UniformBuffer_InvalidArrayStride_Struct) {
// struct ArrayElem {
// a : f32;
// b : i32;
// }
// type Inner = array<ArrayElem, 10>;
//
// [[block]]
// struct Outer {
// inner : Inner;
// scalar : i32;
// };
//
// @group(0) @binding(0)
// var<uniform> a : Outer;
auto* array_elem = Structure("ArrayElem", {
Member("a", ty.f32()),
Member("b", ty.i32()),
});
Alias("Inner", ty.array(ty.Of(array_elem), 10));
Structure(Source{{12, 34}}, "Outer",
{
Member("inner", ty.type_name(Source{{34, 56}}, "Inner")),
Member("scalar", ty.i32()),
},
{StructBlock()});
Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
ast::StorageClass::kUniform, GroupAndBinding(0, 0));
ASSERT_FALSE(r()->Resolve());
EXPECT_EQ(
r()->error(),
R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 8. Consider using the @size attribute on the last struct member.
12:34 note: see layout of struct: 12:34 note: see layout of struct:
/* align(4) size(84) */ struct Outer { /* align(4) size(84) */ struct Outer {
/* offset( 0) align(4) size(80) */ inner : @stride(8) array<f32, 10>; /* offset( 0) align(4) size(80) */ inner : array<ArrayElem, 10>;
/* offset(80) align(4) size( 4) */ scalar : i32; /* offset(80) align(4) size( 4) */ scalar : i32;
/* */ }; /* */ };
78:90 note: see declaration of variable)"); 78:90 note: see declaration of variable)");

4
test/array/assign_to_function_var.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -23,7 +23,7 @@ fn foo(src_param : ArrayType) {
var dst : ArrayType; var dst : ArrayType;
// Assign from type constructor. // Assign from type constructor.
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
// Assign from parameter. // Assign from parameter.
dst = src_param; dst = src_param;

35
test/array/assign_to_function_var.wgsl.expected.hlsl
View File

@ -3,23 +3,20 @@ void unused_entry_point() {
return; return;
} }
struct tint_padded_array_element {
int el;
};
struct S { struct S {
tint_padded_array_element arr[4]; int4 arr[4];
}; };
static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; static int4 src_private[4] = (int4[4])0;
groupshared tint_padded_array_element src_workgroup[4]; groupshared int4 src_workgroup[4];
cbuffer cbuffer_src_uniform : register(b0, space0) { cbuffer cbuffer_src_uniform : register(b0, space0) {
uint4 src_uniform[4]; uint4 src_uniform[4];
}; };
RWByteAddressBuffer src_storage : register(u1, space0); RWByteAddressBuffer src_storage : register(u1, space0);
typedef tint_padded_array_element ret_arr_ret[4]; typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() { ret_arr_ret ret_arr() {
const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5; return tint_symbol_5;
} }
@ -28,37 +25,37 @@ S ret_struct_arr() {
return tint_symbol_6; return tint_symbol_6;
} }
typedef tint_padded_array_element tint_symbol_1_ret[4]; typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) { tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0; int4 arr_1[4] = (int4[4])0;
{ {
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) { [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
const uint scalar_offset = ((offset + (i * 16u))) / 4; const uint scalar_offset = ((offset + (i * 16u))) / 4;
arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]); arr_1[i] = asint(buffer[scalar_offset / 4]);
} }
} }
return arr_1; return arr_1;
} }
typedef tint_padded_array_element tint_symbol_3_ret[4]; typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) { tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0; int4 arr_2[4] = (int4[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u)))); arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
} }
} }
return arr_2; return arr_2;
} }
void foo(tint_padded_array_element src_param[4]) { void foo(int4 src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; int4 src_function[4] = (int4[4])0;
tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0; int4 tint_symbol[4] = (int4[4])0;
const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7; tint_symbol = tint_symbol_7;
tint_symbol = src_param; tint_symbol = src_param;
tint_symbol = ret_arr(); tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0; const int4 src_let[4] = (int4[4])0;
tint_symbol = src_let; tint_symbol = src_let;
tint_symbol = src_function; tint_symbol = src_function;
tint_symbol = src_private; tint_symbol = src_private;

8
test/array/assign_to_function_var.wgsl.expected.msl
View File

@ -1,12 +1,8 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
struct tint_padded_array_element {
/* 0x0000 */ int el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ int4 arr[4];
}; };
struct S { struct S {
/* 0x0000 */ tint_array_wrapper arr; /* 0x0000 */ tint_array_wrapper arr;
@ -34,7 +30,7 @@ S ret_struct_arr() {
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) { void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
tint_array_wrapper src_function = {}; tint_array_wrapper src_function = {};
tint_array_wrapper dst = {}; tint_array_wrapper dst = {};
tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}}; tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
dst = tint_symbol_2; dst = tint_symbol_2;
dst = src_param; dst = src_param;
dst = ret_arr(); dst = ret_arr();

110
test/array/assign_to_function_var.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 60 ; Bound: 64
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -22,7 +22,7 @@
OpName %dst "dst" OpName %dst "dst"
OpName %dst_nested "dst_nested" OpName %dst_nested "dst_nested"
OpName %src_nested "src_nested" OpName %src_nested "src_nested"
OpDecorate %_arr_int_uint_4 ArrayStride 16 OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %S Block OpDecorate %S Block
OpMemberDecorate %S 0 Offset 0 OpMemberDecorate %S 0 Offset 0
OpDecorate %src_uniform NonWritable OpDecorate %src_uniform NonWritable
@ -34,80 +34,84 @@
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4 %_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 %_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
%7 = OpConstantNull %_arr_int_uint_4 %8 = OpConstantNull %_arr_v4int_uint_4
%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 %src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 %_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup %src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_int_uint_4 %S = OpTypeStruct %_arr_v4int_uint_4
%_ptr_Uniform_S = OpTypePointer Uniform %S %_ptr_Uniform_S = OpTypePointer Uniform %S
%src_uniform = OpVariable %_ptr_Uniform_S Uniform %src_uniform = OpVariable %_ptr_Uniform_S Uniform
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
%void = OpTypeVoid %void = OpTypeVoid
%15 = OpTypeFunction %void %16 = OpTypeFunction %void
%19 = OpTypeFunction %_arr_int_uint_4 %20 = OpTypeFunction %_arr_v4int_uint_4
%22 = OpTypeFunction %S %23 = OpTypeFunction %S
%25 = OpConstantNull %S %26 = OpConstantNull %S
%26 = OpTypeFunction %void %_arr_int_uint_4 %27 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 %_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%35 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2 %int_2 = OpConstant %int 2
%37 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3 %int_3 = OpConstant %int 3
%36 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3 %39 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
%40 = OpConstantComposite %_arr_v4int_uint_4 %35 %37 %39 %39
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 %_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 %_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%_arr_int_uint_2 = OpTypeArray %int %uint_2 %_arr_int_uint_2 = OpTypeArray %int %uint_2
%uint_3 = OpConstant %uint 3 %uint_3 = OpConstant %uint 3
%_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3 %_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
%_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4 %_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
%57 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 %61 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
%unused_entry_point = OpFunction %void None %15 %unused_entry_point = OpFunction %void None %16
%18 = OpLabel %19 = OpLabel
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%ret_arr = OpFunction %_arr_int_uint_4 None %19 %ret_arr = OpFunction %_arr_v4int_uint_4 None %20
%21 = OpLabel %22 = OpLabel
OpReturnValue %7 OpReturnValue %8
OpFunctionEnd OpFunctionEnd
%ret_struct_arr = OpFunction %S None %22 %ret_struct_arr = OpFunction %S None %23
%24 = OpLabel %25 = OpLabel
OpReturnValue %25 OpReturnValue %26
OpFunctionEnd OpFunctionEnd
%foo = OpFunction %void None %26 %foo = OpFunction %void None %27
%src_param = OpFunctionParameter %_arr_int_uint_4 %src_param = OpFunctionParameter %_arr_v4int_uint_4
%29 = OpLabel %30 = OpLabel
%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 %src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
%dst = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 %dst = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
%dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %57 %dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
%src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %57 %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
OpStore %dst %36
OpStore %dst %src_param
%37 = OpFunctionCall %_arr_int_uint_4 %ret_arr
OpStore %dst %37
OpStore %dst %7
%38 = OpLoad %_arr_int_uint_4 %src_function
OpStore %dst %38
%39 = OpLoad %_arr_int_uint_4 %src_private
OpStore %dst %39
%40 = OpLoad %_arr_int_uint_4 %src_workgroup
OpStore %dst %40 OpStore %dst %40
%41 = OpFunctionCall %S %ret_struct_arr OpStore %dst %src_param
%42 = OpCompositeExtract %_arr_int_uint_4 %41 0 %41 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
OpStore %dst %41
OpStore %dst %8
%42 = OpLoad %_arr_v4int_uint_4 %src_function
OpStore %dst %42 OpStore %dst %42
%45 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 %43 = OpLoad %_arr_v4int_uint_4 %src_private
%46 = OpLoad %_arr_int_uint_4 %45 OpStore %dst %43
%44 = OpLoad %_arr_v4int_uint_4 %src_workgroup
OpStore %dst %44
%45 = OpFunctionCall %S %ret_struct_arr
%46 = OpCompositeExtract %_arr_v4int_uint_4 %45 0
OpStore %dst %46 OpStore %dst %46
%48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 %49 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
%49 = OpLoad %_arr_int_uint_4 %48 %50 = OpLoad %_arr_v4int_uint_4 %49
OpStore %dst %49 OpStore %dst %50
%59 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
OpStore %dst_nested %59 %53 = OpLoad %_arr_v4int_uint_4 %52
OpStore %dst %53
%63 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
OpStore %dst_nested %63
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

4
test/array/assign_to_function_var.wgsl.expected.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -23,7 +23,7 @@ fn ret_struct_arr() -> S {
fn foo(src_param : ArrayType) { fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
var dst : ArrayType; var dst : ArrayType;
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
dst = src_param; dst = src_param;
dst = ret_arr(); dst = ret_arr();
let src_let : ArrayType = ArrayType(); let src_let : ArrayType = ArrayType();

4
test/array/assign_to_private_var.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -24,7 +24,7 @@ fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
// Assign from type constructor. // Assign from type constructor.
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
// Assign from parameter. // Assign from parameter.
dst = src_param; dst = src_param;

35
test/array/assign_to_private_var.wgsl.expected.hlsl
View File

@ -3,25 +3,22 @@ void unused_entry_point() {
return; return;
} }
struct tint_padded_array_element {
int el;
};
struct S { struct S {
tint_padded_array_element arr[4]; int4 arr[4];
}; };
static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; static int4 src_private[4] = (int4[4])0;
groupshared tint_padded_array_element src_workgroup[4]; groupshared int4 src_workgroup[4];
cbuffer cbuffer_src_uniform : register(b0, space0) { cbuffer cbuffer_src_uniform : register(b0, space0) {
uint4 src_uniform[4]; uint4 src_uniform[4];
}; };
RWByteAddressBuffer src_storage : register(u1, space0); RWByteAddressBuffer src_storage : register(u1, space0);
static tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0; static int4 tint_symbol[4] = (int4[4])0;
static int dst_nested[4][3][2] = (int[4][3][2])0; static int dst_nested[4][3][2] = (int[4][3][2])0;
typedef tint_padded_array_element ret_arr_ret[4]; typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() { ret_arr_ret ret_arr() {
const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5; return tint_symbol_5;
} }
@ -30,36 +27,36 @@ S ret_struct_arr() {
return tint_symbol_6; return tint_symbol_6;
} }
typedef tint_padded_array_element tint_symbol_1_ret[4]; typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) { tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0; int4 arr_1[4] = (int4[4])0;
{ {
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) { [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
const uint scalar_offset = ((offset + (i * 16u))) / 4; const uint scalar_offset = ((offset + (i * 16u))) / 4;
arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]); arr_1[i] = asint(buffer[scalar_offset / 4]);
} }
} }
return arr_1; return arr_1;
} }
typedef tint_padded_array_element tint_symbol_3_ret[4]; typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) { tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0; int4 arr_2[4] = (int4[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u)))); arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
} }
} }
return arr_2; return arr_2;
} }
void foo(tint_padded_array_element src_param[4]) { void foo(int4 src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; int4 src_function[4] = (int4[4])0;
const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7; tint_symbol = tint_symbol_7;
tint_symbol = src_param; tint_symbol = src_param;
tint_symbol = ret_arr(); tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0; const int4 src_let[4] = (int4[4])0;
tint_symbol = src_let; tint_symbol = src_let;
tint_symbol = src_function; tint_symbol = src_function;
tint_symbol = src_private; tint_symbol = src_private;

8
test/array/assign_to_private_var.wgsl.expected.msl
View File

@ -1,12 +1,8 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
struct tint_padded_array_element {
/* 0x0000 */ int el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ int4 arr[4];
}; };
struct S { struct S {
/* 0x0000 */ tint_array_wrapper arr; /* 0x0000 */ tint_array_wrapper arr;
@ -33,7 +29,7 @@ S ret_struct_arr() {
void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, thread tint_array_wrapper_1* const tint_symbol_8) { void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, thread tint_array_wrapper_1* const tint_symbol_8) {
tint_array_wrapper src_function = {}; tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}}; tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
*(tint_symbol_3) = tint_symbol_2; *(tint_symbol_3) = tint_symbol_2;
*(tint_symbol_3) = src_param; *(tint_symbol_3) = src_param;
*(tint_symbol_3) = ret_arr(); *(tint_symbol_3) = ret_arr();

110
test/array/assign_to_private_var.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 61 ; Bound: 65
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -22,7 +22,7 @@
OpName %src_param "src_param" OpName %src_param "src_param"
OpName %src_function "src_function" OpName %src_function "src_function"
OpName %src_nested "src_nested" OpName %src_nested "src_nested"
OpDecorate %_arr_int_uint_4 ArrayStride 16 OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %S Block OpDecorate %S Block
OpMemberDecorate %S 0 Offset 0 OpMemberDecorate %S 0 Offset 0
OpDecorate %src_uniform NonWritable OpDecorate %src_uniform NonWritable
@ -34,81 +34,85 @@
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4 %_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 %_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
%7 = OpConstantNull %_arr_int_uint_4 %8 = OpConstantNull %_arr_v4int_uint_4
%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 %src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 %_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup %src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_int_uint_4 %S = OpTypeStruct %_arr_v4int_uint_4
%_ptr_Uniform_S = OpTypePointer Uniform %S %_ptr_Uniform_S = OpTypePointer Uniform %S
%src_uniform = OpVariable %_ptr_Uniform_S Uniform %src_uniform = OpVariable %_ptr_Uniform_S Uniform
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
%dst = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 %dst = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%_arr_int_uint_2 = OpTypeArray %int %uint_2 %_arr_int_uint_2 = OpTypeArray %int %uint_2
%uint_3 = OpConstant %uint 3 %uint_3 = OpConstant %uint 3
%_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3 %_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
%_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4 %_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
%_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Private %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Private %_arr__arr__arr_int_uint_2_uint_3_uint_4
%23 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 %24 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
%dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %23 %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %24
%void = OpTypeVoid %void = OpTypeVoid
%24 = OpTypeFunction %void %25 = OpTypeFunction %void
%28 = OpTypeFunction %_arr_int_uint_4 %29 = OpTypeFunction %_arr_v4int_uint_4
%31 = OpTypeFunction %S %32 = OpTypeFunction %S
%34 = OpConstantNull %S %35 = OpConstantNull %S
%35 = OpTypeFunction %void %_arr_int_uint_4 %36 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 %_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%43 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2 %int_2 = OpConstant %int 2
%45 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3 %int_3 = OpConstant %int 3
%44 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3 %47 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
%48 = OpConstantComposite %_arr_v4int_uint_4 %43 %45 %47 %47
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 %_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 %_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
%unused_entry_point = OpFunction %void None %24 %unused_entry_point = OpFunction %void None %25
%27 = OpLabel %28 = OpLabel
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%ret_arr = OpFunction %_arr_int_uint_4 None %28 %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
%30 = OpLabel %31 = OpLabel
OpReturnValue %7 OpReturnValue %8
OpFunctionEnd OpFunctionEnd
%ret_struct_arr = OpFunction %S None %31 %ret_struct_arr = OpFunction %S None %32
%33 = OpLabel %34 = OpLabel
OpReturnValue %34 OpReturnValue %35
OpFunctionEnd OpFunctionEnd
%foo = OpFunction %void None %35 %foo = OpFunction %void None %36
%src_param = OpFunctionParameter %_arr_int_uint_4 %src_param = OpFunctionParameter %_arr_v4int_uint_4
%38 = OpLabel %39 = OpLabel
%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 %src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
%src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %23 %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %24
OpStore %dst %44
OpStore %dst %src_param
%45 = OpFunctionCall %_arr_int_uint_4 %ret_arr
OpStore %dst %45
OpStore %dst %7
%46 = OpLoad %_arr_int_uint_4 %src_function
OpStore %dst %46
%47 = OpLoad %_arr_int_uint_4 %src_private
OpStore %dst %47
%48 = OpLoad %_arr_int_uint_4 %src_workgroup
OpStore %dst %48 OpStore %dst %48
%49 = OpFunctionCall %S %ret_struct_arr OpStore %dst %src_param
%50 = OpCompositeExtract %_arr_int_uint_4 %49 0 %49 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
OpStore %dst %49
OpStore %dst %8
%50 = OpLoad %_arr_v4int_uint_4 %src_function
OpStore %dst %50 OpStore %dst %50
%53 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 %51 = OpLoad %_arr_v4int_uint_4 %src_private
%54 = OpLoad %_arr_int_uint_4 %53 OpStore %dst %51
%52 = OpLoad %_arr_v4int_uint_4 %src_workgroup
OpStore %dst %52
%53 = OpFunctionCall %S %ret_struct_arr
%54 = OpCompositeExtract %_arr_v4int_uint_4 %53 0
OpStore %dst %54 OpStore %dst %54
%56 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 %57 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
%57 = OpLoad %_arr_int_uint_4 %56 %58 = OpLoad %_arr_v4int_uint_4 %57
OpStore %dst %57 OpStore %dst %58
%60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
OpStore %dst_nested %60 %61 = OpLoad %_arr_v4int_uint_4 %60
OpStore %dst %61
%64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
OpStore %dst_nested %64
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

4
test/array/assign_to_private_var.wgsl.expected.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -26,7 +26,7 @@ fn ret_struct_arr() -> S {
fn foo(src_param : ArrayType) { fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
dst = src_param; dst = src_param;
dst = ret_arr(); dst = ret_arr();
let src_let : ArrayType = ArrayType(); let src_let : ArrayType = ArrayType();

4
test/array/assign_to_storage_var.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -28,7 +28,7 @@ fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
// Assign from type constructor. // Assign from type constructor.
dst.arr = ArrayType(1, 2, 3, 3); dst.arr = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
// Assign from parameter. // Assign from parameter.
dst.arr = src_param; dst.arr = src_param;

39
test/array/assign_to_storage_var.wgsl.expected.hlsl
View File

@ -3,15 +3,12 @@ void unused_entry_point() {
return; return;
} }
struct tint_padded_array_element {
int el;
};
struct S { struct S {
tint_padded_array_element arr[4]; int4 arr[4];
}; };
static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; static int4 src_private[4] = (int4[4])0;
groupshared tint_padded_array_element src_workgroup[4]; groupshared int4 src_workgroup[4];
cbuffer cbuffer_src_uniform : register(b0, space0) { cbuffer cbuffer_src_uniform : register(b0, space0) {
uint4 src_uniform[4]; uint4 src_uniform[4];
}; };
@ -19,9 +16,9 @@ RWByteAddressBuffer src_storage : register(u1, space0);
RWByteAddressBuffer tint_symbol : register(u2, space0); RWByteAddressBuffer tint_symbol : register(u2, space0);
RWByteAddressBuffer dst_nested : register(u3, space0); RWByteAddressBuffer dst_nested : register(u3, space0);
typedef tint_padded_array_element ret_arr_ret[4]; typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() { ret_arr_ret ret_arr() {
const tint_padded_array_element tint_symbol_11[4] = (tint_padded_array_element[4])0; const int4 tint_symbol_11[4] = (int4[4])0;
return tint_symbol_11; return tint_symbol_11;
} }
@ -30,33 +27,33 @@ S ret_struct_arr() {
return tint_symbol_12; return tint_symbol_12;
} }
void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
tint_padded_array_element array[4] = value; int4 array[4] = value;
{ {
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) { [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
buffer.Store((offset + (i * 16u)), asuint(array[i].el)); buffer.Store4((offset + (i * 16u)), asuint(array[i]));
} }
} }
} }
typedef tint_padded_array_element tint_symbol_3_ret[4]; typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) { tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0; int4 arr_1[4] = (int4[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
const uint scalar_offset = ((offset + (i_1 * 16u))) / 4; const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
arr_1[i_1].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]); arr_1[i_1] = asint(buffer[scalar_offset / 4]);
} }
} }
return arr_1; return arr_1;
} }
typedef tint_padded_array_element tint_symbol_5_ret[4]; typedef int4 tint_symbol_5_ret[4];
tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) { tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0; int4 arr_2[4] = (int4[4])0;
{ {
[loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) { [loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
arr_2[i_2].el = asint(buffer.Load((offset + (i_2 * 16u)))); arr_2[i_2] = asint(buffer.Load4((offset + (i_2 * 16u))));
} }
} }
return arr_2; return arr_2;
@ -89,13 +86,13 @@ void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, int value[4][3][2])
} }
} }
void foo(tint_padded_array_element src_param[4]) { void foo(int4 src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; int4 src_function[4] = (int4[4])0;
const tint_padded_array_element tint_symbol_13[4] = {{1}, {2}, {3}, {3}}; const int4 tint_symbol_13[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol_1(tint_symbol, 0u, tint_symbol_13); tint_symbol_1(tint_symbol, 0u, tint_symbol_13);
tint_symbol_1(tint_symbol, 0u, src_param); tint_symbol_1(tint_symbol, 0u, src_param);
tint_symbol_1(tint_symbol, 0u, ret_arr()); tint_symbol_1(tint_symbol, 0u, ret_arr());
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0; const int4 src_let[4] = (int4[4])0;
tint_symbol_1(tint_symbol, 0u, src_let); tint_symbol_1(tint_symbol, 0u, src_let);
tint_symbol_1(tint_symbol, 0u, src_function); tint_symbol_1(tint_symbol, 0u, src_function);
tint_symbol_1(tint_symbol, 0u, src_private); tint_symbol_1(tint_symbol, 0u, src_private);

8
test/array/assign_to_storage_var.wgsl.expected.msl
View File

@ -1,12 +1,8 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
struct tint_padded_array_element {
/* 0x0000 */ int el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ int4 arr[4];
}; };
struct S { struct S {
/* 0x0000 */ tint_array_wrapper arr; /* 0x0000 */ tint_array_wrapper arr;
@ -36,7 +32,7 @@ S ret_struct_arr() {
void foo(tint_array_wrapper src_param, device S* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) { void foo(tint_array_wrapper src_param, device S* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) {
tint_array_wrapper src_function = {}; tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}}; tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
(*(tint_symbol_3)).arr = tint_symbol_2; (*(tint_symbol_3)).arr = tint_symbol_2;
(*(tint_symbol_3)).arr = src_param; (*(tint_symbol_3)).arr = src_param;
(*(tint_symbol_3)).arr = ret_arr(); (*(tint_symbol_3)).arr = ret_arr();

132
test/array/assign_to_storage_var.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 74 ; Bound: 78
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -24,7 +24,7 @@
OpName %src_param "src_param" OpName %src_param "src_param"
OpName %src_function "src_function" OpName %src_function "src_function"
OpName %src_nested "src_nested" OpName %src_nested "src_nested"
OpDecorate %_arr_int_uint_4 ArrayStride 16 OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %S Block OpDecorate %S Block
OpMemberDecorate %S 0 Offset 0 OpMemberDecorate %S 0 Offset 0
OpDecorate %src_uniform NonWritable OpDecorate %src_uniform NonWritable
@ -42,15 +42,16 @@
OpDecorate %dst_nested DescriptorSet 0 OpDecorate %dst_nested DescriptorSet 0
OpDecorate %dst_nested Binding 3 OpDecorate %dst_nested Binding 3
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4 %_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 %_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
%7 = OpConstantNull %_arr_int_uint_4 %8 = OpConstantNull %_arr_v4int_uint_4
%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 %src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 %_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup %src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_int_uint_4 %S = OpTypeStruct %_arr_v4int_uint_4
%_ptr_Uniform_S = OpTypePointer Uniform %S %_ptr_Uniform_S = OpTypePointer Uniform %S
%src_uniform = OpVariable %_ptr_Uniform_S Uniform %src_uniform = OpVariable %_ptr_Uniform_S Uniform
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
@ -65,71 +66,74 @@
%_ptr_StorageBuffer_S_nested = OpTypePointer StorageBuffer %S_nested %_ptr_StorageBuffer_S_nested = OpTypePointer StorageBuffer %S_nested
%dst_nested = OpVariable %_ptr_StorageBuffer_S_nested StorageBuffer %dst_nested = OpVariable %_ptr_StorageBuffer_S_nested StorageBuffer
%void = OpTypeVoid %void = OpTypeVoid
%24 = OpTypeFunction %void %25 = OpTypeFunction %void
%28 = OpTypeFunction %_arr_int_uint_4 %29 = OpTypeFunction %_arr_v4int_uint_4
%31 = OpTypeFunction %S %32 = OpTypeFunction %S
%34 = OpConstantNull %S %35 = OpConstantNull %S
%35 = OpTypeFunction %void %_arr_int_uint_4 %36 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 %_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 %_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%46 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2 %int_2 = OpConstant %int 2
%48 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3 %int_3 = OpConstant %int 3
%47 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3 %50 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 %51 = OpConstantComposite %_arr_v4int_uint_4 %46 %48 %50 %50
%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
%70 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 %74 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
%_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_int_uint_2_uint_3_uint_4
%unused_entry_point = OpFunction %void None %24 %unused_entry_point = OpFunction %void None %25
%27 = OpLabel %28 = OpLabel
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%ret_arr = OpFunction %_arr_int_uint_4 None %28 %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
%30 = OpLabel %31 = OpLabel
OpReturnValue %7 OpReturnValue %8
OpFunctionEnd OpFunctionEnd
%ret_struct_arr = OpFunction %S None %31 %ret_struct_arr = OpFunction %S None %32
%33 = OpLabel %34 = OpLabel
OpReturnValue %34 OpReturnValue %35
OpFunctionEnd OpFunctionEnd
%foo = OpFunction %void None %35 %foo = OpFunction %void None %36
%src_param = OpFunctionParameter %_arr_int_uint_4 %src_param = OpFunctionParameter %_arr_v4int_uint_4
%38 = OpLabel %39 = OpLabel
%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 %src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
%src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %70 %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %74
%43 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %44 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %43 %47 OpStore %44 %51
%48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %48 %src_param OpStore %52 %src_param
%49 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%50 = OpFunctionCall %_arr_int_uint_4 %ret_arr %54 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
OpStore %49 %50 OpStore %53 %54
%51 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %51 %7 OpStore %55 %8
%52 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%53 = OpLoad %_arr_int_uint_4 %src_function %57 = OpLoad %_arr_v4int_uint_4 %src_function
OpStore %52 %53
%54 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
%55 = OpLoad %_arr_int_uint_4 %src_private
OpStore %54 %55
%56 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
%57 = OpLoad %_arr_int_uint_4 %src_workgroup
OpStore %56 %57 OpStore %56 %57
%58 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%59 = OpFunctionCall %S %ret_struct_arr %59 = OpLoad %_arr_v4int_uint_4 %src_private
%60 = OpCompositeExtract %_arr_int_uint_4 %59 0 OpStore %58 %59
OpStore %58 %60 %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%61 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %61 = OpLoad %_arr_v4int_uint_4 %src_workgroup
%63 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 OpStore %60 %61
%64 = OpLoad %_arr_int_uint_4 %63 %62 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %61 %64 %63 = OpFunctionCall %S %ret_struct_arr
%65 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 %64 = OpCompositeExtract %_arr_v4int_uint_4 %63 0
%66 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 OpStore %62 %64
%67 = OpLoad %_arr_int_uint_4 %66 %65 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
OpStore %65 %67 %67 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
%72 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0 %68 = OpLoad %_arr_v4int_uint_4 %67
%73 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested OpStore %65 %68
OpStore %72 %73 %69 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
%70 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
%71 = OpLoad %_arr_v4int_uint_4 %70
OpStore %69 %71
%76 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0
%77 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
OpStore %76 %77
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

4
test/array/assign_to_storage_var.wgsl.expected.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -30,7 +30,7 @@ fn ret_struct_arr() -> S {
fn foo(src_param : ArrayType) { fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
dst.arr = ArrayType(1, 2, 3, 3); dst.arr = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
dst.arr = src_param; dst.arr = src_param;
dst.arr = ret_arr(); dst.arr = ret_arr();
let src_let : ArrayType = ArrayType(); let src_let : ArrayType = ArrayType();

4
test/array/assign_to_workgroup_var.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -24,7 +24,7 @@ fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
// Assign from type constructor. // Assign from type constructor.
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
// Assign from parameter. // Assign from parameter.
dst = src_param; dst = src_param;

35
test/array/assign_to_workgroup_var.wgsl.expected.hlsl
View File

@ -3,25 +3,22 @@ void unused_entry_point() {
return; return;
} }
struct tint_padded_array_element {
int el;
};
struct S { struct S {
tint_padded_array_element arr[4]; int4 arr[4];
}; };
static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0; static int4 src_private[4] = (int4[4])0;
groupshared tint_padded_array_element src_workgroup[4]; groupshared int4 src_workgroup[4];
cbuffer cbuffer_src_uniform : register(b0, space0) { cbuffer cbuffer_src_uniform : register(b0, space0) {
uint4 src_uniform[4]; uint4 src_uniform[4];
}; };
RWByteAddressBuffer src_storage : register(u1, space0); RWByteAddressBuffer src_storage : register(u1, space0);
groupshared tint_padded_array_element tint_symbol[4]; groupshared int4 tint_symbol[4];
groupshared int dst_nested[4][3][2]; groupshared int dst_nested[4][3][2];
typedef tint_padded_array_element ret_arr_ret[4]; typedef int4 ret_arr_ret[4];
ret_arr_ret ret_arr() { ret_arr_ret ret_arr() {
const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0; const int4 tint_symbol_5[4] = (int4[4])0;
return tint_symbol_5; return tint_symbol_5;
} }
@ -30,36 +27,36 @@ S ret_struct_arr() {
return tint_symbol_6; return tint_symbol_6;
} }
typedef tint_padded_array_element tint_symbol_1_ret[4]; typedef int4 tint_symbol_1_ret[4];
tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) { tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0; int4 arr_1[4] = (int4[4])0;
{ {
[loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) { [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
const uint scalar_offset = ((offset + (i * 16u))) / 4; const uint scalar_offset = ((offset + (i * 16u))) / 4;
arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]); arr_1[i] = asint(buffer[scalar_offset / 4]);
} }
} }
return arr_1; return arr_1;
} }
typedef tint_padded_array_element tint_symbol_3_ret[4]; typedef int4 tint_symbol_3_ret[4];
tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) { tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0; int4 arr_2[4] = (int4[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u)))); arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
} }
} }
return arr_2; return arr_2;
} }
void foo(tint_padded_array_element src_param[4]) { void foo(int4 src_param[4]) {
tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; int4 src_function[4] = (int4[4])0;
const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
tint_symbol = tint_symbol_7; tint_symbol = tint_symbol_7;
tint_symbol = src_param; tint_symbol = src_param;
tint_symbol = ret_arr(); tint_symbol = ret_arr();
const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0; const int4 src_let[4] = (int4[4])0;
tint_symbol = src_let; tint_symbol = src_let;
tint_symbol = src_function; tint_symbol = src_function;
tint_symbol = src_private; tint_symbol = src_private;

8
test/array/assign_to_workgroup_var.wgsl.expected.msl
View File

@ -1,12 +1,8 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
struct tint_padded_array_element {
/* 0x0000 */ int el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ int4 arr[4];
}; };
struct S { struct S {
/* 0x0000 */ tint_array_wrapper arr; /* 0x0000 */ tint_array_wrapper arr;
@ -33,7 +29,7 @@ S ret_struct_arr() {
void foo(tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array_wrapper_1* const tint_symbol_8) { void foo(tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array_wrapper_1* const tint_symbol_8) {
tint_array_wrapper src_function = {}; tint_array_wrapper src_function = {};
tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}}; tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
*(tint_symbol_3) = tint_symbol_2; *(tint_symbol_3) = tint_symbol_2;
*(tint_symbol_3) = src_param; *(tint_symbol_3) = src_param;
*(tint_symbol_3) = ret_arr(); *(tint_symbol_3) = ret_arr();

108
test/array/assign_to_workgroup_var.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 61 ; Bound: 65
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -22,7 +22,7 @@
OpName %src_param "src_param" OpName %src_param "src_param"
OpName %src_function "src_function" OpName %src_function "src_function"
OpName %src_nested "src_nested" OpName %src_nested "src_nested"
OpDecorate %_arr_int_uint_4 ArrayStride 16 OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpDecorate %S Block OpDecorate %S Block
OpMemberDecorate %S 0 Offset 0 OpMemberDecorate %S 0 Offset 0
OpDecorate %src_uniform NonWritable OpDecorate %src_uniform NonWritable
@ -34,20 +34,21 @@
OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4 %_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 %_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
%7 = OpConstantNull %_arr_int_uint_4 %8 = OpConstantNull %_arr_v4int_uint_4
%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 %src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 %_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup %src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%S = OpTypeStruct %_arr_int_uint_4 %S = OpTypeStruct %_arr_v4int_uint_4
%_ptr_Uniform_S = OpTypePointer Uniform %S %_ptr_Uniform_S = OpTypePointer Uniform %S
%src_uniform = OpVariable %_ptr_Uniform_S Uniform %src_uniform = OpVariable %_ptr_Uniform_S Uniform
%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
%src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
%dst = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup %dst = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%_arr_int_uint_2 = OpTypeArray %int %uint_2 %_arr_int_uint_2 = OpTypeArray %int %uint_2
%uint_3 = OpConstant %uint 3 %uint_3 = OpConstant %uint 3
@ -56,59 +57,62 @@
%_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Workgroup %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Workgroup %_arr__arr__arr_int_uint_2_uint_3_uint_4
%dst_nested = OpVariable %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 Workgroup %dst_nested = OpVariable %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 Workgroup
%void = OpTypeVoid %void = OpTypeVoid
%23 = OpTypeFunction %void %24 = OpTypeFunction %void
%27 = OpTypeFunction %_arr_int_uint_4 %28 = OpTypeFunction %_arr_v4int_uint_4
%30 = OpTypeFunction %S %31 = OpTypeFunction %S
%33 = OpConstantNull %S %34 = OpConstantNull %S
%34 = OpTypeFunction %void %_arr_int_uint_4 %35 = OpTypeFunction %void %_arr_v4int_uint_4
%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 %_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%42 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
%int_2 = OpConstant %int 2 %int_2 = OpConstant %int 2
%44 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
%int_3 = OpConstant %int 3 %int_3 = OpConstant %int 3
%43 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3 %46 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
%47 = OpConstantComposite %_arr_v4int_uint_4 %42 %44 %46 %46
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 %_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 %_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
%_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
%59 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 %63 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
%unused_entry_point = OpFunction %void None %23 %unused_entry_point = OpFunction %void None %24
%26 = OpLabel %27 = OpLabel
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%ret_arr = OpFunction %_arr_int_uint_4 None %27 %ret_arr = OpFunction %_arr_v4int_uint_4 None %28
%29 = OpLabel %30 = OpLabel
OpReturnValue %7 OpReturnValue %8
OpFunctionEnd OpFunctionEnd
%ret_struct_arr = OpFunction %S None %30 %ret_struct_arr = OpFunction %S None %31
%32 = OpLabel %33 = OpLabel
OpReturnValue %33 OpReturnValue %34
OpFunctionEnd OpFunctionEnd
%foo = OpFunction %void None %34 %foo = OpFunction %void None %35
%src_param = OpFunctionParameter %_arr_int_uint_4 %src_param = OpFunctionParameter %_arr_v4int_uint_4
%37 = OpLabel %38 = OpLabel
%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 %src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
%src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %59 %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %63
OpStore %dst %43
OpStore %dst %src_param
%44 = OpFunctionCall %_arr_int_uint_4 %ret_arr
OpStore %dst %44
OpStore %dst %7
%45 = OpLoad %_arr_int_uint_4 %src_function
OpStore %dst %45
%46 = OpLoad %_arr_int_uint_4 %src_private
OpStore %dst %46
%47 = OpLoad %_arr_int_uint_4 %src_workgroup
OpStore %dst %47 OpStore %dst %47
%48 = OpFunctionCall %S %ret_struct_arr OpStore %dst %src_param
%49 = OpCompositeExtract %_arr_int_uint_4 %48 0 %48 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
OpStore %dst %48
OpStore %dst %8
%49 = OpLoad %_arr_v4int_uint_4 %src_function
OpStore %dst %49 OpStore %dst %49
%52 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 %50 = OpLoad %_arr_v4int_uint_4 %src_private
%53 = OpLoad %_arr_int_uint_4 %52 OpStore %dst %50
%51 = OpLoad %_arr_v4int_uint_4 %src_workgroup
OpStore %dst %51
%52 = OpFunctionCall %S %ret_struct_arr
%53 = OpCompositeExtract %_arr_v4int_uint_4 %52 0
OpStore %dst %53 OpStore %dst %53
%55 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 %56 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
%56 = OpLoad %_arr_int_uint_4 %55 %57 = OpLoad %_arr_v4int_uint_4 %56
OpStore %dst %56 OpStore %dst %57
%60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested %59 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
OpStore %dst_nested %60 %60 = OpLoad %_arr_v4int_uint_4 %59
OpStore %dst %60
%64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
OpStore %dst_nested %64
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

4
test/array/assign_to_workgroup_var.wgsl.expected.wgsl
View File

@ -1,4 +1,4 @@
type ArrayType = @stride(16) array<i32, 4>; type ArrayType = array<vec4<i32>, 4>;
struct S { struct S {
arr : ArrayType; arr : ArrayType;
@ -26,7 +26,7 @@ fn ret_struct_arr() -> S {
fn foo(src_param : ArrayType) { fn foo(src_param : ArrayType) {
var src_function : ArrayType; var src_function : ArrayType;
dst = ArrayType(1, 2, 3, 3); dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
dst = src_param; dst = src_param;
dst = ret_arr(); dst = ret_arr();
let src_let : ArrayType = ArrayType(); let src_let : ArrayType = ArrayType();

2
test/buffer/storage/dynamic_index/read.wgsl
View File

@ -7,7 +7,7 @@ struct Inner {
f : f32; f : f32;
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : @stride(16) array<vec4<i32>, 4>; i : array<vec4<i32>, 4>;
}; };
struct S { struct S {

2
test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
View File

@ -7,7 +7,7 @@ struct Inner {
f : f32; f : f32;
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : @stride(16) array<vec4<i32>, 4>; i : array<vec4<i32>, 4>;
} }
struct S { struct S {

4
test/buffer/storage/dynamic_index/write.wgsl
View File

@ -7,7 +7,7 @@ struct Inner {
f : f32; f : f32;
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : @stride(16) array<vec4<i32>, 4>; i : array<vec4<i32>, 4>;
}; };
struct S { struct S {
@ -26,5 +26,5 @@ fn main(@builtin(local_invocation_index) idx : u32) {
s.arr[idx].f = f32(); s.arr[idx].f = f32();
s.arr[idx].g = mat2x3<f32>(); s.arr[idx].g = mat2x3<f32>();
s.arr[idx].h = mat3x2<f32>(); s.arr[idx].h = mat3x2<f32>();
s.arr[idx].i = @stride(16) array<vec4<i32>, 4>(); s.arr[idx].i = array<vec4<i32>, 4>();
} }

4
test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
View File

@ -7,7 +7,7 @@ struct Inner {
f : f32; f : f32;
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : @stride(16) array<vec4<i32>, 4>; i : array<vec4<i32>, 4>;
} }
struct S { struct S {
@ -26,5 +26,5 @@ fn main(@builtin(local_invocation_index) idx : u32) {
s.arr[idx].f = f32(); s.arr[idx].f = f32();
s.arr[idx].g = mat2x3<f32>(); s.arr[idx].g = mat2x3<f32>();
s.arr[idx].h = mat3x2<f32>(); s.arr[idx].h = mat3x2<f32>();
s.arr[idx].i = @stride(16) array<vec4<i32>, 4>(); s.arr[idx].i = array<vec4<i32>, 4>();
} }

2
test/buffer/storage/static_index/read.wgsl
View File

@ -12,7 +12,7 @@ struct S {
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : Inner; i : Inner;
j : @stride(16) array<Inner, 4>; j : array<Inner, 4>;
}; };
@binding(0) @group(0) var<storage, read> s : S; @binding(0) @group(0) var<storage, read> s : S;

9
test/buffer/storage/static_index/read.wgsl.expected.glsl
View File

@ -4,9 +4,6 @@ precision mediump float;
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
struct S { struct S {
ivec3 a; ivec3 a;
int b; int b;
@ -17,7 +14,7 @@ struct S {
mat2x3 g; mat2x3 g;
mat3x2 h; mat3x2 h;
Inner i; Inner i;
tint_padded_array_element j[4]; Inner j[4];
}; };
layout (binding = 0) buffer S_1 { layout (binding = 0) buffer S_1 {
@ -30,7 +27,7 @@ layout (binding = 0) buffer S_1 {
mat2x3 g; mat2x3 g;
mat3x2 h; mat3x2 h;
Inner i; Inner i;
tint_padded_array_element j[4]; Inner j[4];
} s; } s;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@ -44,7 +41,7 @@ void tint_symbol() {
mat2x3 g = s.g; mat2x3 g = s.g;
mat3x2 h = s.h; mat3x2 h = s.h;
Inner i = s.i; Inner i = s.i;
tint_padded_array_element j[4] = s.j; Inner j[4] = s.j;
return; return;
} }
void main() { void main() {

11
test/buffer/storage/static_index/read.wgsl.expected.hlsl
View File

@ -1,9 +1,6 @@
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
ByteAddressBuffer s : register(t0, space0); ByteAddressBuffer s : register(t0, space0);
@ -20,12 +17,12 @@ Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
return tint_symbol_11; return tint_symbol_11;
} }
typedef tint_padded_array_element tint_symbol_10_ret[4]; typedef Inner tint_symbol_10_ret[4];
tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) { tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
tint_padded_array_element arr[4] = (tint_padded_array_element[4])0; Inner arr[4] = (Inner[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1].el = tint_symbol_9(buffer, (offset + (i_1 * 16u))); arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 4u)));
} }
} }
return arr; return arr;
@ -42,6 +39,6 @@ void main() {
const float2x3 g = tint_symbol_6(s, 48u); const float2x3 g = tint_symbol_6(s, 48u);
const float3x2 h = tint_symbol_7(s, 80u); const float3x2 h = tint_symbol_7(s, 80u);
const Inner i = tint_symbol_9(s, 104u); const Inner i = tint_symbol_9(s, 104u);
const tint_padded_array_element j[4] = tint_symbol_10(s, 108u); const Inner j[4] = tint_symbol_10(s, 108u);
return; return;
} }

8
test/buffer/storage/static_index/read.wgsl.expected.msl
View File

@ -15,12 +15,8 @@ inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
struct Inner { struct Inner {
/* 0x0000 */ int x; /* 0x0000 */ int x;
}; };
struct tint_padded_array_element {
/* 0x0000 */ Inner el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ Inner arr[4];
}; };
struct S { struct S {
/* 0x0000 */ packed_int3 a; /* 0x0000 */ packed_int3 a;
@ -33,7 +29,7 @@ struct S {
/* 0x0050 */ float3x2 h; /* 0x0050 */ float3x2 h;
/* 0x0068 */ Inner i; /* 0x0068 */ Inner i;
/* 0x006c */ tint_array_wrapper j; /* 0x006c */ tint_array_wrapper j;
/* 0x00ac */ int8_t tint_pad_1[4]; /* 0x007c */ int8_t tint_pad[4];
}; };
kernel void tint_symbol(const device S* tint_symbol_1 [[buffer(0)]]) { kernel void tint_symbol(const device S* tint_symbol_1 [[buffer(0)]]) {

2
test/buffer/storage/static_index/read.wgsl.expected.spvasm
View File

@ -38,7 +38,7 @@
OpMemberDecorate %S 8 Offset 104 OpMemberDecorate %S 8 Offset 104
OpMemberDecorate %Inner 0 Offset 0 OpMemberDecorate %Inner 0 Offset 0
OpMemberDecorate %S 9 Offset 108 OpMemberDecorate %S 9 Offset 108
OpDecorate %_arr_Inner_uint_4 ArrayStride 16 OpDecorate %_arr_Inner_uint_4 ArrayStride 4
OpDecorate %s NonWritable OpDecorate %s NonWritable
OpDecorate %s Binding 0 OpDecorate %s Binding 0
OpDecorate %s DescriptorSet 0 OpDecorate %s DescriptorSet 0

2
test/buffer/storage/static_index/read.wgsl.expected.wgsl
View File

@ -12,7 +12,7 @@ struct S {
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : Inner; i : Inner;
j : @stride(16) array<Inner, 4>; j : array<Inner, 4>;
} }
@binding(0) @group(0) var<storage, read> s : S; @binding(0) @group(0) var<storage, read> s : S;

4
test/buffer/storage/static_index/write.wgsl
View File

@ -12,7 +12,7 @@ struct S {
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : Inner; i : Inner;
j : @stride(16) array<Inner, 4>; j : array<Inner, 4>;
}; };
@binding(0) @group(0) var<storage, write> s : S; @binding(0) @group(0) var<storage, write> s : S;
@ -28,5 +28,5 @@ fn main() {
s.g = mat2x3<f32>(); s.g = mat2x3<f32>();
s.h = mat3x2<f32>(); s.h = mat3x2<f32>();
s.i = Inner(); s.i = Inner();
s.j = @stride(16) array<Inner, 4>(); s.j = array<Inner, 4>();
} }

9
test/buffer/storage/static_index/write.wgsl.expected.glsl
View File

@ -4,9 +4,6 @@ precision mediump float;
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
struct S { struct S {
ivec3 a; ivec3 a;
int b; int b;
@ -17,7 +14,7 @@ struct S {
mat2x3 g; mat2x3 g;
mat3x2 h; mat3x2 h;
Inner i; Inner i;
tint_padded_array_element j[4]; Inner j[4];
}; };
layout (binding = 0) buffer S_1 { layout (binding = 0) buffer S_1 {
@ -30,7 +27,7 @@ layout (binding = 0) buffer S_1 {
mat2x3 g; mat2x3 g;
mat3x2 h; mat3x2 h;
Inner i; Inner i;
tint_padded_array_element j[4]; Inner j[4];
} s; } s;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@ -45,7 +42,7 @@ void tint_symbol() {
s.h = mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f); s.h = mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
Inner tint_symbol_1 = Inner(0); Inner tint_symbol_1 = Inner(0);
s.i = tint_symbol_1; s.i = tint_symbol_1;
tint_padded_array_element tint_symbol_2[4] = tint_padded_array_element[4](tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0))); Inner tint_symbol_2[4] = Inner[4](Inner(0), Inner(0), Inner(0), Inner(0));
s.j = tint_symbol_2; s.j = tint_symbol_2;
return; return;
} }

11
test/buffer/storage/static_index/write.wgsl.expected.hlsl
View File

@ -1,9 +1,6 @@
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
RWByteAddressBuffer s : register(u0, space0); RWByteAddressBuffer s : register(u0, space0);
@ -22,11 +19,11 @@ void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, Inner value) {
buffer.Store((offset + 0u), asuint(value.x)); buffer.Store((offset + 0u), asuint(value.x));
} }
void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, Inner value[4]) {
tint_padded_array_element array[4] = value; Inner array[4] = value;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1].el); tint_symbol_9(buffer, (offset + (i_1 * 4u)), array[i_1]);
} }
} }
} }
@ -43,7 +40,7 @@ void main() {
tint_symbol_7(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f)); tint_symbol_7(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
const Inner tint_symbol_11 = (Inner)0; const Inner tint_symbol_11 = (Inner)0;
tint_symbol_9(s, 104u, tint_symbol_11); tint_symbol_9(s, 104u, tint_symbol_11);
const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0; const Inner tint_symbol_12[4] = (Inner[4])0;
tint_symbol_10(s, 108u, tint_symbol_12); tint_symbol_10(s, 108u, tint_symbol_12);
return; return;
} }

8
test/buffer/storage/static_index/write.wgsl.expected.msl
View File

@ -15,12 +15,8 @@ inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
struct Inner { struct Inner {
/* 0x0000 */ int x; /* 0x0000 */ int x;
}; };
struct tint_padded_array_element {
/* 0x0000 */ Inner el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ Inner arr[4];
}; };
struct S { struct S {
/* 0x0000 */ packed_int3 a; /* 0x0000 */ packed_int3 a;
@ -33,7 +29,7 @@ struct S {
/* 0x0050 */ float3x2 h; /* 0x0050 */ float3x2 h;
/* 0x0068 */ Inner i; /* 0x0068 */ Inner i;
/* 0x006c */ tint_array_wrapper j; /* 0x006c */ tint_array_wrapper j;
/* 0x00ac */ int8_t tint_pad_1[4]; /* 0x007c */ int8_t tint_pad[4];
}; };
kernel void tint_symbol(device S* tint_symbol_3 [[buffer(0)]]) { kernel void tint_symbol(device S* tint_symbol_3 [[buffer(0)]]) {

2
test/buffer/storage/static_index/write.wgsl.expected.spvasm
View File

@ -38,7 +38,7 @@
OpMemberDecorate %S 8 Offset 104 OpMemberDecorate %S 8 Offset 104
OpMemberDecorate %Inner 0 Offset 0 OpMemberDecorate %Inner 0 Offset 0
OpMemberDecorate %S 9 Offset 108 OpMemberDecorate %S 9 Offset 108
OpDecorate %_arr_Inner_uint_4 ArrayStride 16 OpDecorate %_arr_Inner_uint_4 ArrayStride 4
OpDecorate %s NonReadable OpDecorate %s NonReadable
OpDecorate %s Binding 0 OpDecorate %s Binding 0
OpDecorate %s DescriptorSet 0 OpDecorate %s DescriptorSet 0

4
test/buffer/storage/static_index/write.wgsl.expected.wgsl
View File

@ -12,7 +12,7 @@ struct S {
g : mat2x3<f32>; g : mat2x3<f32>;
h : mat3x2<f32>; h : mat3x2<f32>;
i : Inner; i : Inner;
j : @stride(16) array<Inner, 4>; j : array<Inner, 4>;
} }
@binding(0) @group(0) var<storage, write> s : S; @binding(0) @group(0) var<storage, write> s : S;
@ -28,5 +28,5 @@ fn main() {
s.g = mat2x3<f32>(); s.g = mat2x3<f32>();
s.h = mat3x2<f32>(); s.h = mat3x2<f32>();
s.i = Inner(); s.i = Inner();
s.j = @stride(16) array<Inner, 4>(); s.j = array<Inner, 4>();
} }

2
test/buffer/uniform/dynamic_index/read.wgsl
View File

@ -9,7 +9,7 @@ struct Inner {
h : vec2<i32>; h : vec2<i32>;
i : mat2x3<f32>; i : mat2x3<f32>;
@align(16) j : mat3x2<f32>; @align(16) j : mat3x2<f32>;
@align(16) k : @stride(16) array<vec4<i32>, 4>; @align(16) k : array<vec4<i32>, 4>;
}; };
struct S { struct S {

2
test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
View File

@ -11,7 +11,7 @@ struct Inner {
@align(16) @align(16)
j : mat3x2<f32>; j : mat3x2<f32>;
@align(16) @align(16)
k : @stride(16) array<vec4<i32>, 4>; k : array<vec4<i32>, 4>;
} }
struct S { struct S {

4
test/buffer/uniform/static_index/read.wgsl
View File

@ -1,5 +1,5 @@
struct Inner { struct Inner {
x : i32; @size(16) x : i32;
}; };
struct S { struct S {
@ -14,7 +14,7 @@ struct S {
i : mat2x3<f32>; i : mat2x3<f32>;
j : mat3x2<f32>; j : mat3x2<f32>;
@align(16) k : Inner; @align(16) k : Inner;
@align(16) l : @stride(16) array<Inner, 4>; @align(16) l : array<Inner, 4>;
}; };
@binding(0) @group(0) var<uniform> s : S; @binding(0) @group(0) var<uniform> s : S;

9
test/buffer/uniform/static_index/read.wgsl.expected.glsl
View File

@ -4,9 +4,6 @@ precision mediump float;
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
struct S { struct S {
ivec3 a; ivec3 a;
int b; int b;
@ -19,7 +16,7 @@ struct S {
mat2x3 i; mat2x3 i;
mat3x2 j; mat3x2 j;
Inner k; Inner k;
tint_padded_array_element l[4]; Inner l[4];
}; };
layout (binding = 0) uniform S_1 { layout (binding = 0) uniform S_1 {
@ -34,7 +31,7 @@ layout (binding = 0) uniform S_1 {
mat2x3 i; mat2x3 i;
mat3x2 j; mat3x2 j;
Inner k; Inner k;
tint_padded_array_element l[4]; Inner l[4];
} s; } s;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@ -50,7 +47,7 @@ void tint_symbol() {
mat2x3 i = s.i; mat2x3 i = s.i;
mat3x2 j = s.j; mat3x2 j = s.j;
Inner k = s.k; Inner k = s.k;
tint_padded_array_element l[4] = s.l; Inner l[4] = s.l;
return; return;
} }
void main() { void main() {

11
test/buffer/uniform/static_index/read.wgsl.expected.hlsl
View File

@ -1,9 +1,6 @@
struct Inner { struct Inner {
int x; int x;
}; };
struct tint_padded_array_element {
Inner el;
};
cbuffer cbuffer_s : register(b0, space0) { cbuffer cbuffer_s : register(b0, space0) {
uint4 s[13]; uint4 s[13];
@ -31,12 +28,12 @@ Inner tint_symbol_10(uint4 buffer[13], uint offset) {
return tint_symbol_12; return tint_symbol_12;
} }
typedef tint_padded_array_element tint_symbol_11_ret[4]; typedef Inner tint_symbol_11_ret[4];
tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) { tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
tint_padded_array_element arr[4] = (tint_padded_array_element[4])0; Inner arr[4] = (Inner[4])0;
{ {
[loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) { [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
arr[i_1].el = tint_symbol_10(buffer, (offset + (i_1 * 16u))); arr[i_1] = tint_symbol_10(buffer, (offset + (i_1 * 16u)));
} }
} }
return arr; return arr;
@ -55,6 +52,6 @@ void main() {
const float2x3 i = tint_symbol_7(s, 64u); const float2x3 i = tint_symbol_7(s, 64u);
const float3x2 j = tint_symbol_8(s, 96u); const float3x2 j = tint_symbol_8(s, 96u);
const Inner k = tint_symbol_10(s, 128u); const Inner k = tint_symbol_10(s, 128u);
const tint_padded_array_element l[4] = tint_symbol_11(s, 144u); const Inner l[4] = tint_symbol_11(s, 144u);
return; return;
} }

6
test/buffer/uniform/static_index/read.wgsl.expected.msl
View File

@ -14,13 +14,10 @@ inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
struct Inner { struct Inner {
/* 0x0000 */ int x; /* 0x0000 */ int x;
};
struct tint_padded_array_element {
/* 0x0000 */ Inner el;
/* 0x0004 */ int8_t tint_pad[12]; /* 0x0004 */ int8_t tint_pad[12];
}; };
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ Inner arr[4];
}; };
struct S { struct S {
/* 0x0000 */ packed_int3 a; /* 0x0000 */ packed_int3 a;
@ -35,7 +32,6 @@ struct S {
/* 0x0060 */ float3x2 j; /* 0x0060 */ float3x2 j;
/* 0x0078 */ int8_t tint_pad_1[8]; /* 0x0078 */ int8_t tint_pad_1[8];
/* 0x0080 */ Inner k; /* 0x0080 */ Inner k;
/* 0x0084 */ int8_t tint_pad_2[12];
/* 0x0090 */ tint_array_wrapper l; /* 0x0090 */ tint_array_wrapper l;
}; };

3
test/buffer/uniform/static_index/read.wgsl.expected.wgsl
View File

@ -1,4 +1,5 @@
struct Inner { struct Inner {
@size(16)
x : i32; x : i32;
} }
@ -16,7 +17,7 @@ struct S {
@align(16) @align(16)
k : Inner; k : Inner;
@align(16) @align(16)
l : @stride(16) array<Inner, 4>; l : array<Inner, 4>;
} }
@binding(0) @group(0) var<uniform> s : S; @binding(0) @group(0) var<uniform> s : S;

10
test/bug/chromium/1273230.wgsl
View File

@ -37,23 +37,23 @@ struct Dbg {
}; };
struct F32s { struct F32s {
values : @stride(4) array<f32>; values : array<f32>;
}; };
struct U32s { struct U32s {
values : @stride(4) array<u32>; values : array<u32>;
}; };
struct I32s { struct I32s {
values : @stride(4) array<i32>; values : array<i32>;
}; };
struct AU32s { struct AU32s {
values : @stride(4) array<atomic<u32>>; values : array<atomic<u32>>;
}; };
struct AI32s { struct AI32s {
values : @stride(4) array<atomic<i32>>; values : array<atomic<i32>>;
}; };
@binding(0) @group(0) var<uniform> uniforms : Uniforms; @binding(0) @group(0) var<uniform> uniforms : Uniforms;

10
test/bug/chromium/1273230.wgsl.expected.wgsl
View File

@ -54,23 +54,23 @@ struct Dbg {
} }
struct F32s { struct F32s {
values : @stride(4) array<f32>; values : array<f32>;
} }
struct U32s { struct U32s {
values : @stride(4) array<u32>; values : array<u32>;
} }
struct I32s { struct I32s {
values : @stride(4) array<i32>; values : array<i32>;
} }
struct AU32s { struct AU32s {
values : @stride(4) array<atomic<u32>>; values : array<atomic<u32>>;
} }
struct AI32s { struct AI32s {
values : @stride(4) array<atomic<i32>>; values : array<atomic<i32>>;
} }
@binding(0) @group(0) var<uniform> uniforms : Uniforms; @binding(0) @group(0) var<uniform> uniforms : Uniforms;

4
test/bug/fxc/dyn_array_idx/read/uniform.wgsl
View File

@ -1,5 +1,5 @@
struct UBO { struct UBO {
data: @stride(16) array<i32, 4>; data: array<vec4<i32>, 4>;
dynamic_idx: i32; dynamic_idx: i32;
}; };
@group(0) @binding(0) var<uniform> ubo: UBO; @group(0) @binding(0) var<uniform> ubo: UBO;
@ -10,5 +10,5 @@ struct Result {
@stage(compute) @workgroup_size(1) @stage(compute) @workgroup_size(1)
fn f() { fn f() {
result.out = ubo.data[ubo.dynamic_idx]; result.out = ubo.data[ubo.dynamic_idx].x;
} }

9
test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl
View File

@ -1,16 +1,13 @@
#version 310 es #version 310 es
precision mediump float; precision mediump float;
struct tint_padded_array_element {
int el;
};
struct UBO { struct UBO {
tint_padded_array_element data[4]; ivec4 data[4];
int dynamic_idx; int dynamic_idx;
}; };
layout (binding = 0) uniform UBO_1 { layout (binding = 0) uniform UBO_1 {
tint_padded_array_element data[4]; ivec4 data[4];
int dynamic_idx; int dynamic_idx;
} ubo; } ubo;
@ -24,7 +21,7 @@ layout (binding = 2) buffer Result_1 {
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void f() { void f() {
result.tint_symbol = ubo.data[ubo.dynamic_idx].el; result.tint_symbol = ubo.data[ubo.dynamic_idx].x;
return; return;
} }
void main() { void main() {

9
test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
View File

@ -1,23 +1,20 @@
#include <metal_stdlib> #include <metal_stdlib>
using namespace metal; using namespace metal;
struct tint_padded_array_element {
/* 0x0000 */ int el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper { struct tint_array_wrapper {
/* 0x0000 */ tint_padded_array_element arr[4]; /* 0x0000 */ int4 arr[4];
}; };
struct UBO { struct UBO {
/* 0x0000 */ tint_array_wrapper data; /* 0x0000 */ tint_array_wrapper data;
/* 0x0040 */ int dynamic_idx; /* 0x0040 */ int dynamic_idx;
/* 0x0044 */ int8_t tint_pad[12];
}; };
struct Result { struct Result {
/* 0x0000 */ int out; /* 0x0000 */ int out;
}; };
kernel void f(device Result* tint_symbol [[buffer(1)]], const constant UBO* tint_symbol_1 [[buffer(0)]]) { kernel void f(device Result* tint_symbol [[buffer(1)]], const constant UBO* tint_symbol_1 [[buffer(0)]]) {
(*(tint_symbol)).out = (*(tint_symbol_1)).data.arr[(*(tint_symbol_1)).dynamic_idx].el; (*(tint_symbol)).out = (*(tint_symbol_1)).data.arr[(*(tint_symbol_1)).dynamic_idx][0];
return; return;
} }

27
test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.3 ; Version: 1.3
; Generator: Google Tint Compiler; 0 ; Generator: Google Tint Compiler; 0
; Bound: 24 ; Bound: 25
; Schema: 0 ; Schema: 0
OpCapability Shader OpCapability Shader
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
@ -17,7 +17,7 @@
OpName %f "f" OpName %f "f"
OpDecorate %UBO Block OpDecorate %UBO Block
OpMemberDecorate %UBO 0 Offset 0 OpMemberDecorate %UBO 0 Offset 0
OpDecorate %_arr_int_uint_4 ArrayStride 16 OpDecorate %_arr_v4int_uint_4 ArrayStride 16
OpMemberDecorate %UBO 1 Offset 64 OpMemberDecorate %UBO 1 Offset 64
OpDecorate %ubo NonWritable OpDecorate %ubo NonWritable
OpDecorate %ubo DescriptorSet 0 OpDecorate %ubo DescriptorSet 0
@ -27,28 +27,29 @@
OpDecorate %result DescriptorSet 0 OpDecorate %result DescriptorSet 0
OpDecorate %result Binding 2 OpDecorate %result Binding 2
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%v4int = OpTypeVector %int 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_4 = OpConstant %uint 4 %uint_4 = OpConstant %uint 4
%_arr_int_uint_4 = OpTypeArray %int %uint_4 %_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
%UBO = OpTypeStruct %_arr_int_uint_4 %int %UBO = OpTypeStruct %_arr_v4int_uint_4 %int
%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO %_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
%ubo = OpVariable %_ptr_Uniform_UBO Uniform %ubo = OpVariable %_ptr_Uniform_UBO Uniform
%Result = OpTypeStruct %int %Result = OpTypeStruct %int
%_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result %_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
%result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
%void = OpTypeVoid %void = OpTypeVoid
%11 = OpTypeFunction %void %12 = OpTypeFunction %void
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%_ptr_Uniform_int = OpTypePointer Uniform %int %_ptr_Uniform_int = OpTypePointer Uniform %int
%f = OpFunction %void None %11 %f = OpFunction %void None %12
%14 = OpLabel %15 = OpLabel
%17 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0 %18 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
%20 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1 %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1
%21 = OpLoad %int %20 %22 = OpLoad %int %21
%22 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %21 %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %22 %uint_0
%23 = OpLoad %int %22 %24 = OpLoad %int %23
OpStore %17 %23 OpStore %18 %24
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

4
test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl
View File

@ -1,5 +1,5 @@
struct UBO { struct UBO {
data : @stride(16) array<i32, 4>; data : array<vec4<i32>, 4>;
dynamic_idx : i32; dynamic_idx : i32;
} }
@ -13,5 +13,5 @@ struct Result {
@stage(compute) @workgroup_size(1) @stage(compute) @workgroup_size(1)
fn f() { fn f() {
result.out = ubo.data[ubo.dynamic_idx]; result.out = ubo.data[ubo.dynamic_idx].x;
} }

2
test/bug/tint/1046.wgsl
View File

@ -4,7 +4,7 @@ struct PointLight {
}; };
struct PointLights { struct PointLights {
values : @stride(16) array<PointLight>; values : array<PointLight>;
}; };
struct Uniforms { struct Uniforms {

2
test/bug/tint/1046.wgsl.expected.wgsl
View File

@ -3,7 +3,7 @@ struct PointLight {
} }
struct PointLights { struct PointLights {
values : @stride(16) array<PointLight>; values : array<PointLight>;
} }
struct Uniforms { struct Uniforms {

10
test/bug/tint/1113.wgsl
View File

@ -23,11 +23,11 @@
value_f32_3 : f32; value_f32_3 : f32;
}; };
struct F32s { values : @stride(4) array<f32>; }; struct F32s { values : array<f32>; };
struct U32s { values : @stride(4) array<u32>; }; struct U32s { values : array<u32>; };
struct I32s { values : @stride(4) array<i32>; }; struct I32s { values : array<i32>; };
struct AU32s { values : @stride(4) array<atomic<u32>>; }; struct AU32s { values : array<atomic<u32>>; };
struct AI32s { values : @stride(4) array<atomic<i32>>; }; struct AI32s { values : array<atomic<i32>>; };
// IN // IN
@binding(0) @group(0) var<uniform> uniforms : Uniforms; @binding(0) @group(0) var<uniform> uniforms : Uniforms;

10
test/bug/tint/1113.wgsl.expected.wgsl
View File

@ -23,23 +23,23 @@ struct Dbg {
} }
struct F32s { struct F32s {
values : @stride(4) array<f32>; values : array<f32>;
} }
struct U32s { struct U32s {
values : @stride(4) array<u32>; values : array<u32>;
} }
struct I32s { struct I32s {
values : @stride(4) array<i32>; values : array<i32>;
} }
struct AU32s { struct AU32s {
values : @stride(4) array<atomic<u32>>; values : array<atomic<u32>>;
} }
struct AI32s { struct AI32s {
values : @stride(4) array<atomic<i32>>; values : array<atomic<i32>>;
} }
@binding(0) @group(0) var<uniform> uniforms : Uniforms; @binding(0) @group(0) var<uniform> uniforms : Uniforms;

2
test/bug/tint/294.wgsl
View File

@ -3,6 +3,6 @@ struct Light {
colour : vec3<f32>; colour : vec3<f32>;
}; };
struct Lights { struct Lights {
light : @stride(32) array<Light>; light : array<Light>;
}; };
@group(0) @binding(1) var<storage, read> lights : Lights; @group(0) @binding(1) var<storage, read> lights : Lights;

2
test/bug/tint/294.wgsl.expected.wgsl
View File

@ -4,7 +4,7 @@ struct Light {
} }
struct Lights { struct Lights {
light : @stride(32) array<Light>; light : array<Light>;
} }
@group(0) @binding(1) var<storage, read> lights : Lights; @group(0) @binding(1) var<storage, read> lights : Lights;

2
test/bug/tint/534.wgsl
View File

@ -5,7 +5,7 @@
channelCount : u32; channelCount : u32;
}; };
struct OutputBuf { struct OutputBuf {
result : @stride(4) array<u32>; result : array<u32>;
}; };
@group(0) @binding(0) var src : texture_2d<f32>; @group(0) @binding(0) var src : texture_2d<f32>;
@group(0) @binding(1) var dst : texture_2d<f32>; @group(0) @binding(1) var dst : texture_2d<f32>;

2
test/bug/tint/534.wgsl.expected.wgsl
View File

@ -6,7 +6,7 @@ struct Uniforms {
} }
struct OutputBuf { struct OutputBuf {
result : @stride(4) array<u32>; result : array<u32>;
} }
@group(0) @binding(0) var src : texture_2d<f32>; @group(0) @binding(0) var src : texture_2d<f32>;

2
test/bug/tint/757.wgsl
View File

@ -7,7 +7,7 @@
@group(0) @binding(1) var myTexture : texture_2d_array<f32>; @group(0) @binding(1) var myTexture : texture_2d_array<f32>;
struct Result { struct Result {
values : @stride(4) array<f32>; values : array<f32>;
}; };
@group(0) @binding(3) var<storage, read_write> result : Result; @group(0) @binding(3) var<storage, read_write> result : Result;

2
test/bug/tint/757.wgsl.expected.wgsl
View File

@ -7,7 +7,7 @@ struct Constants {
@group(0) @binding(1) var myTexture : texture_2d_array<f32>; @group(0) @binding(1) var myTexture : texture_2d_array<f32>;
struct Result { struct Result {
values : @stride(4) array<f32>; values : array<f32>;
} }
@group(0) @binding(3) var<storage, read_write> result : Result; @group(0) @binding(3) var<storage, read_write> result : Result;

2
test/bug/tint/782.wgsl
View File

@ -1,4 +1,4 @@
type ArrayExplicitStride = @stride(4) array<i32, 2>; type ArrayExplicitStride = array<i32, 2>;
type ArrayImplicitStride = array<i32, 2>; type ArrayImplicitStride = array<i32, 2>;
fn foo() { fn foo() {

2
test/bug/tint/782.wgsl.expected.wgsl
View File

@ -1,4 +1,4 @@
type ArrayExplicitStride = @stride(4) array<i32, 2>; type ArrayExplicitStride = array<i32, 2>;
type ArrayImplicitStride = array<i32, 2>; type ArrayImplicitStride = array<i32, 2>;

4
test/bug/tint/922.wgsl
View File

@ -21,12 +21,12 @@ struct ub_SceneParams {
}; };
struct ub_MaterialParams { struct ub_MaterialParams {
u_TexMtx: @stride(32) array<Mat4x2_,1>; u_TexMtx: array<Mat4x2_,1>;
u_Misc0_: vec4<f32>; u_Misc0_: vec4<f32>;
}; };
struct ub_PacketParams { struct ub_PacketParams {
u_PosMtx: @stride(48) array<Mat4x3_,32>; u_PosMtx: array<Mat4x3_,32>;
}; };
struct VertexOutput { struct VertexOutput {

4
test/bug/tint/922.wgsl.expected.wgsl
View File

@ -21,12 +21,12 @@ struct ub_SceneParams {
} }
struct ub_MaterialParams { struct ub_MaterialParams {
u_TexMtx : @stride(32) array<Mat4x2_, 1>; u_TexMtx : array<Mat4x2_, 1>;
u_Misc0_ : vec4<f32>; u_Misc0_ : vec4<f32>;
} }
struct ub_PacketParams { struct ub_PacketParams {
u_PosMtx : @stride(48) array<Mat4x3_, 32>; u_PosMtx : array<Mat4x3_, 32>;
} }
struct VertexOutput { struct VertexOutput {