diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md index bc03ac616e..435934ecbf 100644 --- a/docs/origin-trial-changes.md +++ b/docs/origin-trial-changes.md @@ -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) * 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) +* The `[[stride]]` attribute has been deprecated. [tint:1381](https://crbug.com/tint/1381) ### New Features diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc index 9f543090af..6570e431aa 100644 --- a/src/reader/wgsl/parser_impl.cc +++ b/src/reader/wgsl/parser_impl.cc @@ -3076,7 +3076,9 @@ Maybe ParserImpl::decoration() { auto val = expect_nonzero_positive_sint(use); if (val.errored) return Failure::kErrored; - + deprecated(t.source(), + "the @stride attribute is deprecated; use a larger type if " + "necessary"); return create(t.source(), val.value); }); } diff --git a/src/reader/wgsl/parser_impl_error_msg_test.cc b/src/reader/wgsl/parser_impl_error_msg_test.cc index dfb5c57de5..3c7b91d9f6 100644 --- a/src/reader/wgsl/parser_impl_error_msg_test.cc +++ b/src/reader/wgsl/parser_impl_error_msg_test.cc @@ -985,24 +985,24 @@ var i : array;", + "var i : [[location(1) array;", 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; +var i : [[location(1) array; ^^ -test.wgsl:1:21 error: expected ']]' for decoration list -var i : [[stride(1) array; - ^^^^^ +test.wgsl:1:23 error: expected ']]' for decoration list +var i : [[location(1) array; + ^^^^^ )"); } @@ -1025,14 +1025,14 @@ var i : [[stride 1)]] array; TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayDecoStrideMissingRParen) { EXPECT( - "var i : [[stride(1]] array;", + "var i : [[location(1]] array;", 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; +var i : [[location(1]] array; ^^ -test.wgsl:1:19 error: expected ')' for stride decoration -var i : [[stride(1]] array; - ^^ +test.wgsl:1:21 error: expected ')' for location decoration +var i : [[location(1]] array; + ^^ )"); } diff --git a/src/reader/wgsl/parser_impl_type_decl_test.cc b/src/reader/wgsl/parser_impl_type_decl_test.cc index 13e368a1b6..476b19ed22 100644 --- a/src/reader/wgsl/parser_impl_type_decl_test.cc +++ b/src/reader/wgsl/parser_impl_type_decl_test.cc @@ -534,7 +534,10 @@ TEST_F(ParserImplTest, TypeDecl_Array_Decoration_MissingArray) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); 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) { @@ -546,6 +549,7 @@ TEST_F(ParserImplTest, TypeDecl_Array_Decoration_UnknownDecoration) { ASSERT_TRUE(p->has_error()); EXPECT_EQ(p->error(), R"(1:2: expected decoration)"); } + TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) { auto p = parser("@stride 4) array"); auto t = p->type_decl(); @@ -563,7 +567,10 @@ TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingRightParen) { EXPECT_FALSE(t.matched); ASSERT_EQ(t.value, nullptr); 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) { @@ -610,6 +617,7 @@ TEST_F(ParserImplTest, EXPECT_EQ( p->error(), 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)"); } @@ -638,6 +646,7 @@ TEST_F(ParserImplTest, DEPRECATED_TypeDecl_Array_Stride_MissingRightParen) { EXPECT_EQ( p->error(), 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)"); } diff --git a/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc index 73efe8ccd8..3c43ed8175 100644 --- a/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc +++ b/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc @@ -77,7 +77,7 @@ TEST_F(ParserImplTest, VariableIdentDecl_InvalidIdent) { } 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{}); ast::StructMemberList members; @@ -94,11 +94,11 @@ TEST_F(ParserImplTest, VariableIdentDecl_NonAccessDecoFail) { } 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"); ASSERT_TRUE(p->has_error()); 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) { @@ -112,27 +112,27 @@ TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingLeftParen) { // TODO(crbug.com/tint/1382): Remove TEST_F(ParserImplTest, 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"); ASSERT_TRUE(p->has_error()); ASSERT_TRUE(decl.errored); ASSERT_EQ( p->error(), 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 TEST_F(ParserImplTest, 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"); ASSERT_TRUE(p->has_error()); ASSERT_TRUE(decl.errored); ASSERT_EQ( p->error(), 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 diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc index 5edc4ee994..8864a26de1 100644 --- a/src/resolver/resolver_validation.cc +++ b/src/resolver/resolver_validation.cc @@ -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 // 16 bytes. 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(); + vec && vec->type()->Size() == 4) { + hint = "Consider using a vec4 instead."; + } else if (arr->ElemType()->Is()) { + 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( "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 " + - std::to_string(arr->Stride()) + - ". Consider setting @stride(" + - std::to_string( - utils::RoundUp(required_align, arr->Stride())) + - ") on the array type", + std::to_string(arr->Stride()) + ". " + hint, m->Declaration()->type->source); AddNote("see layout of struct:\n" + str->Layout(builder_->Symbols()), str->Declaration()->source); diff --git a/src/resolver/storage_class_layout_validation_test.cc b/src/resolver/storage_class_layout_validation_test.cc index f9fcc121c8..a8e7056c73 100644 --- a/src/resolver/storage_class_layout_validation_test.cc +++ b/src/resolver/storage_class_layout_validation_test.cc @@ -378,8 +378,8 @@ TEST_F(ResolverStorageClassLayoutValidationTest, // Detect array stride must be a multiple of 16 bytes for uniform buffers TEST_F(ResolverStorageClassLayoutValidationTest, - UniformBuffer_InvalidArrayStride) { - // type Inner = @stride(8) array; + UniformBuffer_InvalidArrayStride_Scalar) { + // type Inner = array; // // [[block]] // struct Outer { @@ -390,7 +390,7 @@ TEST_F(ResolverStorageClassLayoutValidationTest, // @group(0) @binding(0) // var a : Outer; - Alias("Inner", ty.array(ty.f32(), 10, 8)); + Alias("Inner", ty.array(ty.f32(), 10)); Structure(Source{{12, 34}}, "Outer", { @@ -405,10 +405,93 @@ TEST_F(ResolverStorageClassLayoutValidationTest, 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 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; +/* offset(40) align(4) size( 4) */ scalar : i32; +/* */ }; +78:90 note: see declaration of variable)"); +} + +TEST_F(ResolverStorageClassLayoutValidationTest, + UniformBuffer_InvalidArrayStride_Vector) { + // type Inner = array, 10>; + // + // [[block]] + // struct Outer { + // inner : Inner; + // scalar : i32; + // }; + // + // @group(0) @binding(0) + // var a : Outer; + + Alias("Inner", ty.array(ty.vec2(), 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, 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; + // + // [[block]] + // struct Outer { + // inner : Inner; + // scalar : i32; + // }; + // + // @group(0) @binding(0) + // var 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: /* align(4) size(84) */ struct Outer { -/* offset( 0) align(4) size(80) */ inner : @stride(8) array; +/* offset( 0) align(4) size(80) */ inner : array; /* offset(80) align(4) size( 4) */ scalar : i32; /* */ }; 78:90 note: see declaration of variable)"); diff --git a/test/array/assign_to_function_var.wgsl b/test/array/assign_to_function_var.wgsl index d5ccaa26f2..b2dc68d8f9 100644 --- a/test/array/assign_to_function_var.wgsl +++ b/test/array/assign_to_function_var.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -23,7 +23,7 @@ fn foo(src_param : ArrayType) { var dst : ArrayType; // Assign from type constructor. - dst = ArrayType(1, 2, 3, 3); + dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3)); // Assign from parameter. dst = src_param; diff --git a/test/array/assign_to_function_var.wgsl.expected.hlsl b/test/array/assign_to_function_var.wgsl.expected.hlsl index f5c4ff0a76..eda01ad8f9 100644 --- a/test/array/assign_to_function_var.wgsl.expected.hlsl +++ b/test/array/assign_to_function_var.wgsl.expected.hlsl @@ -3,23 +3,20 @@ void unused_entry_point() { return; } -struct tint_padded_array_element { - int el; -}; 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; -groupshared tint_padded_array_element src_workgroup[4]; +static int4 src_private[4] = (int4[4])0; +groupshared int4 src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { uint4 src_uniform[4]; }; 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() { - 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; } @@ -28,37 +25,37 @@ S ret_struct_arr() { 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_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)) { 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; } -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_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)) { - 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; } -void foo(tint_padded_array_element src_param[4]) { - tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; - tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0; - const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; +void foo(int4 src_param[4]) { + int4 src_function[4] = (int4[4])0; + int4 tint_symbol[4] = (int4[4])0; + 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 = src_param; 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_function; tint_symbol = src_private; diff --git a/test/array/assign_to_function_var.wgsl.expected.msl b/test/array/assign_to_function_var.wgsl.expected.msl index e198433c61..7abc32da9c 100644 --- a/test/array/assign_to_function_var.wgsl.expected.msl +++ b/test/array/assign_to_function_var.wgsl.expected.msl @@ -1,12 +1,8 @@ #include using namespace metal; -struct tint_padded_array_element { - /* 0x0000 */ int el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ int4 arr[4]; }; struct S { /* 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) { tint_array_wrapper src_function = {}; 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 = src_param; dst = ret_arr(); diff --git a/test/array/assign_to_function_var.wgsl.expected.spvasm b/test/array/assign_to_function_var.wgsl.expected.spvasm index 668fc9c704..636346d438 100644 --- a/test/array/assign_to_function_var.wgsl.expected.spvasm +++ b/test/array/assign_to_function_var.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 60 +; Bound: 64 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -22,7 +22,7 @@ OpName %dst "dst" OpName %dst_nested "dst_nested" OpName %src_nested "src_nested" - OpDecorate %_arr_int_uint_4 ArrayStride 16 + OpDecorate %_arr_v4int_uint_4 ArrayStride 16 OpDecorate %S Block OpMemberDecorate %S 0 Offset 0 OpDecorate %src_uniform NonWritable @@ -34,80 +34,84 @@ OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 %int = OpTypeInt 32 1 + %v4int = OpTypeVector %int 4 %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 -%_arr_int_uint_4 = OpTypeArray %int %uint_4 -%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 - %7 = OpConstantNull %_arr_int_uint_4 -%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 -%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 -%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup - %S = OpTypeStruct %_arr_int_uint_4 +%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4 +%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4 + %8 = OpConstantNull %_arr_v4int_uint_4 +%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8 +%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4 +%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup + %S = OpTypeStruct %_arr_v4int_uint_4 %_ptr_Uniform_S = OpTypePointer Uniform %S %src_uniform = OpVariable %_ptr_Uniform_S Uniform %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer %void = OpTypeVoid - %15 = OpTypeFunction %void - %19 = OpTypeFunction %_arr_int_uint_4 - %22 = OpTypeFunction %S - %25 = OpConstantNull %S - %26 = OpTypeFunction %void %_arr_int_uint_4 -%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 + %16 = OpTypeFunction %void + %20 = OpTypeFunction %_arr_v4int_uint_4 + %23 = OpTypeFunction %S + %26 = OpConstantNull %S + %27 = OpTypeFunction %void %_arr_v4int_uint_4 +%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4 %int_1 = OpConstant %int 1 + %35 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 %int_2 = OpConstant %int 2 + %37 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2 %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 -%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 -%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 +%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4 +%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4 %uint_2 = OpConstant %uint 2 %_arr_int_uint_2 = OpTypeArray %int %uint_2 %uint_3 = OpConstant %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 %_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 -%unused_entry_point = OpFunction %void None %15 - %18 = OpLabel + %61 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 +%unused_entry_point = OpFunction %void None %16 + %19 = OpLabel OpReturn OpFunctionEnd - %ret_arr = OpFunction %_arr_int_uint_4 None %19 - %21 = OpLabel - OpReturnValue %7 + %ret_arr = OpFunction %_arr_v4int_uint_4 None %20 + %22 = OpLabel + OpReturnValue %8 OpFunctionEnd -%ret_struct_arr = OpFunction %S None %22 - %24 = OpLabel - OpReturnValue %25 +%ret_struct_arr = OpFunction %S None %23 + %25 = OpLabel + OpReturnValue %26 OpFunctionEnd - %foo = OpFunction %void None %26 - %src_param = OpFunctionParameter %_arr_int_uint_4 - %29 = OpLabel -%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 - %dst = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 - %dst_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 %57 - 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 + %foo = OpFunction %void None %27 + %src_param = OpFunctionParameter %_arr_v4int_uint_4 + %30 = OpLabel +%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8 + %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 %61 + %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61 OpStore %dst %40 - %41 = OpFunctionCall %S %ret_struct_arr - %42 = OpCompositeExtract %_arr_int_uint_4 %41 0 + OpStore %dst %src_param + %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 - %45 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 - %46 = OpLoad %_arr_int_uint_4 %45 + %43 = OpLoad %_arr_v4int_uint_4 %src_private + 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 - %48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 - %49 = OpLoad %_arr_int_uint_4 %48 - OpStore %dst %49 - %59 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested - OpStore %dst_nested %59 + %49 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 + %50 = OpLoad %_arr_v4int_uint_4 %49 + OpStore %dst %50 + %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 + %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 OpFunctionEnd diff --git a/test/array/assign_to_function_var.wgsl.expected.wgsl b/test/array/assign_to_function_var.wgsl.expected.wgsl index 0033bea13d..b001ab620d 100644 --- a/test/array/assign_to_function_var.wgsl.expected.wgsl +++ b/test/array/assign_to_function_var.wgsl.expected.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -23,7 +23,7 @@ fn ret_struct_arr() -> S { fn foo(src_param : ArrayType) { var src_function : ArrayType; var dst : ArrayType; - dst = ArrayType(1, 2, 3, 3); + dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3)); dst = src_param; dst = ret_arr(); let src_let : ArrayType = ArrayType(); diff --git a/test/array/assign_to_private_var.wgsl b/test/array/assign_to_private_var.wgsl index 510da193e7..c19b3eedb1 100644 --- a/test/array/assign_to_private_var.wgsl +++ b/test/array/assign_to_private_var.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -24,7 +24,7 @@ fn foo(src_param : ArrayType) { var src_function : ArrayType; // Assign from type constructor. - dst = ArrayType(1, 2, 3, 3); + dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3)); // Assign from parameter. dst = src_param; diff --git a/test/array/assign_to_private_var.wgsl.expected.hlsl b/test/array/assign_to_private_var.wgsl.expected.hlsl index 4fd9d453dd..ddb084e623 100644 --- a/test/array/assign_to_private_var.wgsl.expected.hlsl +++ b/test/array/assign_to_private_var.wgsl.expected.hlsl @@ -3,25 +3,22 @@ void unused_entry_point() { return; } -struct tint_padded_array_element { - int el; -}; 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; -groupshared tint_padded_array_element src_workgroup[4]; +static int4 src_private[4] = (int4[4])0; +groupshared int4 src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { uint4 src_uniform[4]; }; 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; -typedef tint_padded_array_element ret_arr_ret[4]; +typedef int4 ret_arr_ret[4]; 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; } @@ -30,36 +27,36 @@ S ret_struct_arr() { 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_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)) { 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; } -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_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)) { - 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; } -void foo(tint_padded_array_element src_param[4]) { - tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; - const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; +void foo(int4 src_param[4]) { + int4 src_function[4] = (int4[4])0; + 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 = src_param; 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_function; tint_symbol = src_private; diff --git a/test/array/assign_to_private_var.wgsl.expected.msl b/test/array/assign_to_private_var.wgsl.expected.msl index dd91dcad76..7d8a3ab764 100644 --- a/test/array/assign_to_private_var.wgsl.expected.msl +++ b/test/array/assign_to_private_var.wgsl.expected.msl @@ -1,12 +1,8 @@ #include using namespace metal; -struct tint_padded_array_element { - /* 0x0000 */ int el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ int4 arr[4]; }; struct S { /* 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) { 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) = src_param; *(tint_symbol_3) = ret_arr(); diff --git a/test/array/assign_to_private_var.wgsl.expected.spvasm b/test/array/assign_to_private_var.wgsl.expected.spvasm index 258a540796..3427ea782f 100644 --- a/test/array/assign_to_private_var.wgsl.expected.spvasm +++ b/test/array/assign_to_private_var.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 61 +; Bound: 65 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -22,7 +22,7 @@ OpName %src_param "src_param" OpName %src_function "src_function" OpName %src_nested "src_nested" - OpDecorate %_arr_int_uint_4 ArrayStride 16 + OpDecorate %_arr_v4int_uint_4 ArrayStride 16 OpDecorate %S Block OpMemberDecorate %S 0 Offset 0 OpDecorate %src_uniform NonWritable @@ -34,81 +34,85 @@ OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 %int = OpTypeInt 32 1 + %v4int = OpTypeVector %int 4 %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 -%_arr_int_uint_4 = OpTypeArray %int %uint_4 -%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 - %7 = OpConstantNull %_arr_int_uint_4 -%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 -%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 -%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup - %S = OpTypeStruct %_arr_int_uint_4 +%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4 +%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4 + %8 = OpConstantNull %_arr_v4int_uint_4 +%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8 +%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4 +%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup + %S = OpTypeStruct %_arr_v4int_uint_4 %_ptr_Uniform_S = OpTypePointer Uniform %S %src_uniform = OpVariable %_ptr_Uniform_S Uniform %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %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 %_arr_int_uint_2 = OpTypeArray %int %uint_2 %uint_3 = OpConstant %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 %_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 - %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %23 + %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 %24 %void = OpTypeVoid - %24 = OpTypeFunction %void - %28 = OpTypeFunction %_arr_int_uint_4 - %31 = OpTypeFunction %S - %34 = OpConstantNull %S - %35 = OpTypeFunction %void %_arr_int_uint_4 -%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 + %25 = OpTypeFunction %void + %29 = OpTypeFunction %_arr_v4int_uint_4 + %32 = OpTypeFunction %S + %35 = OpConstantNull %S + %36 = OpTypeFunction %void %_arr_v4int_uint_4 +%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4 %int_1 = OpConstant %int 1 + %43 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 %int_2 = OpConstant %int 2 + %45 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2 %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 -%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 -%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 +%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_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 -%unused_entry_point = OpFunction %void None %24 - %27 = OpLabel +%unused_entry_point = OpFunction %void None %25 + %28 = OpLabel OpReturn OpFunctionEnd - %ret_arr = OpFunction %_arr_int_uint_4 None %28 - %30 = OpLabel - OpReturnValue %7 + %ret_arr = OpFunction %_arr_v4int_uint_4 None %29 + %31 = OpLabel + OpReturnValue %8 OpFunctionEnd -%ret_struct_arr = OpFunction %S None %31 - %33 = OpLabel - OpReturnValue %34 +%ret_struct_arr = OpFunction %S None %32 + %34 = OpLabel + OpReturnValue %35 OpFunctionEnd - %foo = OpFunction %void None %35 - %src_param = OpFunctionParameter %_arr_int_uint_4 - %38 = OpLabel -%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 - %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %23 - 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 + %foo = OpFunction %void None %36 + %src_param = OpFunctionParameter %_arr_v4int_uint_4 + %39 = OpLabel +%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 %24 OpStore %dst %48 - %49 = OpFunctionCall %S %ret_struct_arr - %50 = OpCompositeExtract %_arr_int_uint_4 %49 0 + OpStore %dst %src_param + %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 - %53 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 - %54 = OpLoad %_arr_int_uint_4 %53 + %51 = OpLoad %_arr_v4int_uint_4 %src_private + 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 - %56 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 - %57 = OpLoad %_arr_int_uint_4 %56 - OpStore %dst %57 - %60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested - OpStore %dst_nested %60 + %57 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 + %58 = OpLoad %_arr_v4int_uint_4 %57 + OpStore %dst %58 + %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 + %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 OpFunctionEnd diff --git a/test/array/assign_to_private_var.wgsl.expected.wgsl b/test/array/assign_to_private_var.wgsl.expected.wgsl index fd2057beee..7f88483b95 100644 --- a/test/array/assign_to_private_var.wgsl.expected.wgsl +++ b/test/array/assign_to_private_var.wgsl.expected.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -26,7 +26,7 @@ fn ret_struct_arr() -> S { fn foo(src_param : 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 = ret_arr(); let src_let : ArrayType = ArrayType(); diff --git a/test/array/assign_to_storage_var.wgsl b/test/array/assign_to_storage_var.wgsl index 6d001d8dc8..9ed30b2cf2 100644 --- a/test/array/assign_to_storage_var.wgsl +++ b/test/array/assign_to_storage_var.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -28,7 +28,7 @@ fn foo(src_param : ArrayType) { var src_function : ArrayType; // 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. dst.arr = src_param; diff --git a/test/array/assign_to_storage_var.wgsl.expected.hlsl b/test/array/assign_to_storage_var.wgsl.expected.hlsl index 8e9d0b6e04..c940b98042 100644 --- a/test/array/assign_to_storage_var.wgsl.expected.hlsl +++ b/test/array/assign_to_storage_var.wgsl.expected.hlsl @@ -3,15 +3,12 @@ void unused_entry_point() { return; } -struct tint_padded_array_element { - int el; -}; 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; -groupshared tint_padded_array_element src_workgroup[4]; +static int4 src_private[4] = (int4[4])0; +groupshared int4 src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { uint4 src_uniform[4]; }; @@ -19,9 +16,9 @@ RWByteAddressBuffer src_storage : register(u1, space0); RWByteAddressBuffer tint_symbol : register(u2, 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() { - 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; } @@ -30,33 +27,33 @@ S ret_struct_arr() { return tint_symbol_12; } -void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { - tint_padded_array_element array[4] = value; +void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, int4 value[4]) { + int4 array[4] = value; { [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_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)) { 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; } -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_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)) { - 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; @@ -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]) { - tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; - const tint_padded_array_element tint_symbol_13[4] = {{1}, {2}, {3}, {3}}; +void foo(int4 src_param[4]) { + int4 src_function[4] = (int4[4])0; + 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, src_param); 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_function); tint_symbol_1(tint_symbol, 0u, src_private); diff --git a/test/array/assign_to_storage_var.wgsl.expected.msl b/test/array/assign_to_storage_var.wgsl.expected.msl index c0d7dba0d8..c1a31f9a2d 100644 --- a/test/array/assign_to_storage_var.wgsl.expected.msl +++ b/test/array/assign_to_storage_var.wgsl.expected.msl @@ -1,12 +1,8 @@ #include using namespace metal; -struct tint_padded_array_element { - /* 0x0000 */ int el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ int4 arr[4]; }; struct S { /* 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) { 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 = src_param; (*(tint_symbol_3)).arr = ret_arr(); diff --git a/test/array/assign_to_storage_var.wgsl.expected.spvasm b/test/array/assign_to_storage_var.wgsl.expected.spvasm index 1764efac33..e4da62afe4 100644 --- a/test/array/assign_to_storage_var.wgsl.expected.spvasm +++ b/test/array/assign_to_storage_var.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 74 +; Bound: 78 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -24,7 +24,7 @@ OpName %src_param "src_param" OpName %src_function "src_function" OpName %src_nested "src_nested" - OpDecorate %_arr_int_uint_4 ArrayStride 16 + OpDecorate %_arr_v4int_uint_4 ArrayStride 16 OpDecorate %S Block OpMemberDecorate %S 0 Offset 0 OpDecorate %src_uniform NonWritable @@ -42,15 +42,16 @@ OpDecorate %dst_nested DescriptorSet 0 OpDecorate %dst_nested Binding 3 %int = OpTypeInt 32 1 + %v4int = OpTypeVector %int 4 %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 -%_arr_int_uint_4 = OpTypeArray %int %uint_4 -%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 - %7 = OpConstantNull %_arr_int_uint_4 -%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 -%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 -%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup - %S = OpTypeStruct %_arr_int_uint_4 +%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4 +%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4 + %8 = OpConstantNull %_arr_v4int_uint_4 +%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8 +%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4 +%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup + %S = OpTypeStruct %_arr_v4int_uint_4 %_ptr_Uniform_S = OpTypePointer Uniform %S %src_uniform = OpVariable %_ptr_Uniform_S Uniform %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S @@ -65,71 +66,74 @@ %_ptr_StorageBuffer_S_nested = OpTypePointer StorageBuffer %S_nested %dst_nested = OpVariable %_ptr_StorageBuffer_S_nested StorageBuffer %void = OpTypeVoid - %24 = OpTypeFunction %void - %28 = OpTypeFunction %_arr_int_uint_4 - %31 = OpTypeFunction %S - %34 = OpConstantNull %S - %35 = OpTypeFunction %void %_arr_int_uint_4 -%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 + %25 = OpTypeFunction %void + %29 = OpTypeFunction %_arr_v4int_uint_4 + %32 = OpTypeFunction %S + %35 = OpConstantNull %S + %36 = OpTypeFunction %void %_arr_v4int_uint_4 +%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4 %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 + %46 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 %int_2 = OpConstant %int 2 + %48 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2 %int_3 = OpConstant %int 3 - %47 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3 -%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 + %50 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3 + %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 - %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 -%unused_entry_point = OpFunction %void None %24 - %27 = OpLabel +%unused_entry_point = OpFunction %void None %25 + %28 = OpLabel OpReturn OpFunctionEnd - %ret_arr = OpFunction %_arr_int_uint_4 None %28 - %30 = OpLabel - OpReturnValue %7 + %ret_arr = OpFunction %_arr_v4int_uint_4 None %29 + %31 = OpLabel + OpReturnValue %8 OpFunctionEnd -%ret_struct_arr = OpFunction %S None %31 - %33 = OpLabel - OpReturnValue %34 +%ret_struct_arr = OpFunction %S None %32 + %34 = OpLabel + OpReturnValue %35 OpFunctionEnd - %foo = OpFunction %void None %35 - %src_param = OpFunctionParameter %_arr_int_uint_4 - %38 = OpLabel -%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 - %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %70 - %43 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - OpStore %43 %47 - %48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - OpStore %48 %src_param - %49 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - %50 = OpFunctionCall %_arr_int_uint_4 %ret_arr - OpStore %49 %50 - %51 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - OpStore %51 %7 - %52 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - %53 = OpLoad %_arr_int_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 + %foo = OpFunction %void None %36 + %src_param = OpFunctionParameter %_arr_v4int_uint_4 + %39 = OpLabel +%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 %74 + %44 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + OpStore %44 %51 + %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + OpStore %52 %src_param + %53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %54 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr + OpStore %53 %54 + %55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + OpStore %55 %8 + %56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %57 = OpLoad %_arr_v4int_uint_4 %src_function OpStore %56 %57 - %58 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - %59 = OpFunctionCall %S %ret_struct_arr - %60 = OpCompositeExtract %_arr_int_uint_4 %59 0 - OpStore %58 %60 - %61 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - %63 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 - %64 = OpLoad %_arr_int_uint_4 %63 - OpStore %61 %64 - %65 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0 - %66 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 - %67 = OpLoad %_arr_int_uint_4 %66 - OpStore %65 %67 - %72 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0 - %73 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested - OpStore %72 %73 + %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %59 = OpLoad %_arr_v4int_uint_4 %src_private + OpStore %58 %59 + %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %61 = OpLoad %_arr_v4int_uint_4 %src_workgroup + OpStore %60 %61 + %62 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %63 = OpFunctionCall %S %ret_struct_arr + %64 = OpCompositeExtract %_arr_v4int_uint_4 %63 0 + OpStore %62 %64 + %65 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0 + %67 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 + %68 = OpLoad %_arr_v4int_uint_4 %67 + OpStore %65 %68 + %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 OpFunctionEnd diff --git a/test/array/assign_to_storage_var.wgsl.expected.wgsl b/test/array/assign_to_storage_var.wgsl.expected.wgsl index 88902fadba..fe474586c7 100644 --- a/test/array/assign_to_storage_var.wgsl.expected.wgsl +++ b/test/array/assign_to_storage_var.wgsl.expected.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -30,7 +30,7 @@ fn ret_struct_arr() -> S { fn foo(src_param : 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 = ret_arr(); let src_let : ArrayType = ArrayType(); diff --git a/test/array/assign_to_workgroup_var.wgsl b/test/array/assign_to_workgroup_var.wgsl index 9fb94eb9dd..16153482be 100644 --- a/test/array/assign_to_workgroup_var.wgsl +++ b/test/array/assign_to_workgroup_var.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -24,7 +24,7 @@ fn foo(src_param : ArrayType) { var src_function : ArrayType; // Assign from type constructor. - dst = ArrayType(1, 2, 3, 3); + dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3)); // Assign from parameter. dst = src_param; diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl index b60b9741ba..4243ec1edb 100644 --- a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl +++ b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl @@ -3,25 +3,22 @@ void unused_entry_point() { return; } -struct tint_padded_array_element { - int el; -}; 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; -groupshared tint_padded_array_element src_workgroup[4]; +static int4 src_private[4] = (int4[4])0; +groupshared int4 src_workgroup[4]; cbuffer cbuffer_src_uniform : register(b0, space0) { uint4 src_uniform[4]; }; 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]; -typedef tint_padded_array_element ret_arr_ret[4]; +typedef int4 ret_arr_ret[4]; 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; } @@ -30,36 +27,36 @@ S ret_struct_arr() { 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_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)) { 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; } -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_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)) { - 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; } -void foo(tint_padded_array_element src_param[4]) { - tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0; - const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}}; +void foo(int4 src_param[4]) { + int4 src_function[4] = (int4[4])0; + 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 = src_param; 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_function; tint_symbol = src_private; diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.msl b/test/array/assign_to_workgroup_var.wgsl.expected.msl index 05115062b9..2c45f26f6e 100644 --- a/test/array/assign_to_workgroup_var.wgsl.expected.msl +++ b/test/array/assign_to_workgroup_var.wgsl.expected.msl @@ -1,12 +1,8 @@ #include using namespace metal; -struct tint_padded_array_element { - /* 0x0000 */ int el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ int4 arr[4]; }; struct S { /* 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) { 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) = src_param; *(tint_symbol_3) = ret_arr(); diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.spvasm b/test/array/assign_to_workgroup_var.wgsl.expected.spvasm index 5f467012e1..a834fb2079 100644 --- a/test/array/assign_to_workgroup_var.wgsl.expected.spvasm +++ b/test/array/assign_to_workgroup_var.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 61 +; Bound: 65 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -22,7 +22,7 @@ OpName %src_param "src_param" OpName %src_function "src_function" OpName %src_nested "src_nested" - OpDecorate %_arr_int_uint_4 ArrayStride 16 + OpDecorate %_arr_v4int_uint_4 ArrayStride 16 OpDecorate %S Block OpMemberDecorate %S 0 Offset 0 OpDecorate %src_uniform NonWritable @@ -34,20 +34,21 @@ OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8 OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24 %int = OpTypeInt 32 1 + %v4int = OpTypeVector %int 4 %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 -%_arr_int_uint_4 = OpTypeArray %int %uint_4 -%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4 - %7 = OpConstantNull %_arr_int_uint_4 -%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7 -%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4 -%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup - %S = OpTypeStruct %_arr_int_uint_4 +%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4 +%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4 + %8 = OpConstantNull %_arr_v4int_uint_4 +%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8 +%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4 +%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup + %S = OpTypeStruct %_arr_v4int_uint_4 %_ptr_Uniform_S = OpTypePointer Uniform %S %src_uniform = OpVariable %_ptr_Uniform_S Uniform %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S %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 %_arr_int_uint_2 = OpTypeArray %int %uint_2 %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 %dst_nested = OpVariable %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 Workgroup %void = OpTypeVoid - %23 = OpTypeFunction %void - %27 = OpTypeFunction %_arr_int_uint_4 - %30 = OpTypeFunction %S - %33 = OpConstantNull %S - %34 = OpTypeFunction %void %_arr_int_uint_4 -%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 + %24 = OpTypeFunction %void + %28 = OpTypeFunction %_arr_v4int_uint_4 + %31 = OpTypeFunction %S + %34 = OpConstantNull %S + %35 = OpTypeFunction %void %_arr_v4int_uint_4 +%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4 %int_1 = OpConstant %int 1 + %42 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1 %int_2 = OpConstant %int 2 + %44 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2 %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 -%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4 -%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4 +%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_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 - %59 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 -%unused_entry_point = OpFunction %void None %23 - %26 = OpLabel + %63 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4 +%unused_entry_point = OpFunction %void None %24 + %27 = OpLabel OpReturn OpFunctionEnd - %ret_arr = OpFunction %_arr_int_uint_4 None %27 - %29 = OpLabel - OpReturnValue %7 + %ret_arr = OpFunction %_arr_v4int_uint_4 None %28 + %30 = OpLabel + OpReturnValue %8 OpFunctionEnd -%ret_struct_arr = OpFunction %S None %30 - %32 = OpLabel - OpReturnValue %33 +%ret_struct_arr = OpFunction %S None %31 + %33 = OpLabel + OpReturnValue %34 OpFunctionEnd - %foo = OpFunction %void None %34 - %src_param = OpFunctionParameter %_arr_int_uint_4 - %37 = OpLabel -%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7 - %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %59 - 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 + %foo = OpFunction %void None %35 + %src_param = OpFunctionParameter %_arr_v4int_uint_4 + %38 = OpLabel +%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 %63 OpStore %dst %47 - %48 = OpFunctionCall %S %ret_struct_arr - %49 = OpCompositeExtract %_arr_int_uint_4 %48 0 + OpStore %dst %src_param + %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 - %52 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0 - %53 = OpLoad %_arr_int_uint_4 %52 + %50 = OpLoad %_arr_v4int_uint_4 %src_private + 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 - %55 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0 - %56 = OpLoad %_arr_int_uint_4 %55 - OpStore %dst %56 - %60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested - OpStore %dst_nested %60 + %56 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0 + %57 = OpLoad %_arr_v4int_uint_4 %56 + OpStore %dst %57 + %59 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0 + %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 OpFunctionEnd diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.wgsl b/test/array/assign_to_workgroup_var.wgsl.expected.wgsl index 4fecbc8115..6888710166 100644 --- a/test/array/assign_to_workgroup_var.wgsl.expected.wgsl +++ b/test/array/assign_to_workgroup_var.wgsl.expected.wgsl @@ -1,4 +1,4 @@ -type ArrayType = @stride(16) array; +type ArrayType = array, 4>; struct S { arr : ArrayType; @@ -26,7 +26,7 @@ fn ret_struct_arr() -> S { fn foo(src_param : 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 = ret_arr(); let src_let : ArrayType = ArrayType(); diff --git a/test/buffer/storage/dynamic_index/read.wgsl b/test/buffer/storage/dynamic_index/read.wgsl index 94c2c1f8ee..5cd0316f55 100644 --- a/test/buffer/storage/dynamic_index/read.wgsl +++ b/test/buffer/storage/dynamic_index/read.wgsl @@ -7,7 +7,7 @@ struct Inner { f : f32; g : mat2x3; h : mat3x2; - i : @stride(16) array, 4>; + i : array, 4>; }; struct S { diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl index fd0c451cf7..f9c274c3a6 100644 --- a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl +++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl @@ -7,7 +7,7 @@ struct Inner { f : f32; g : mat2x3; h : mat3x2; - i : @stride(16) array, 4>; + i : array, 4>; } struct S { diff --git a/test/buffer/storage/dynamic_index/write.wgsl b/test/buffer/storage/dynamic_index/write.wgsl index 3770015432..4167c1e035 100644 --- a/test/buffer/storage/dynamic_index/write.wgsl +++ b/test/buffer/storage/dynamic_index/write.wgsl @@ -7,7 +7,7 @@ struct Inner { f : f32; g : mat2x3; h : mat3x2; - i : @stride(16) array, 4>; + i : array, 4>; }; struct S { @@ -26,5 +26,5 @@ fn main(@builtin(local_invocation_index) idx : u32) { s.arr[idx].f = f32(); s.arr[idx].g = mat2x3(); s.arr[idx].h = mat3x2(); - s.arr[idx].i = @stride(16) array, 4>(); + s.arr[idx].i = array, 4>(); } diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl index d51db2197c..661916d263 100644 --- a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl +++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl @@ -7,7 +7,7 @@ struct Inner { f : f32; g : mat2x3; h : mat3x2; - i : @stride(16) array, 4>; + i : array, 4>; } struct S { @@ -26,5 +26,5 @@ fn main(@builtin(local_invocation_index) idx : u32) { s.arr[idx].f = f32(); s.arr[idx].g = mat2x3(); s.arr[idx].h = mat3x2(); - s.arr[idx].i = @stride(16) array, 4>(); + s.arr[idx].i = array, 4>(); } diff --git a/test/buffer/storage/static_index/read.wgsl b/test/buffer/storage/static_index/read.wgsl index 773fbbb49c..06ea57e9e7 100644 --- a/test/buffer/storage/static_index/read.wgsl +++ b/test/buffer/storage/static_index/read.wgsl @@ -12,7 +12,7 @@ struct S { g : mat2x3; h : mat3x2; i : Inner; - j : @stride(16) array; + j : array; }; @binding(0) @group(0) var s : S; diff --git a/test/buffer/storage/static_index/read.wgsl.expected.glsl b/test/buffer/storage/static_index/read.wgsl.expected.glsl index f481a982eb..1b4c469f25 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.glsl +++ b/test/buffer/storage/static_index/read.wgsl.expected.glsl @@ -4,9 +4,6 @@ precision mediump float; struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; struct S { ivec3 a; int b; @@ -17,7 +14,7 @@ struct S { mat2x3 g; mat3x2 h; Inner i; - tint_padded_array_element j[4]; + Inner j[4]; }; layout (binding = 0) buffer S_1 { @@ -30,7 +27,7 @@ layout (binding = 0) buffer S_1 { mat2x3 g; mat3x2 h; Inner i; - tint_padded_array_element j[4]; + Inner j[4]; } s; 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; mat3x2 h = s.h; Inner i = s.i; - tint_padded_array_element j[4] = s.j; + Inner j[4] = s.j; return; } void main() { diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl index 667739be6a..9062d007ed 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.hlsl +++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl @@ -1,9 +1,6 @@ struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; ByteAddressBuffer s : register(t0, space0); @@ -20,12 +17,12 @@ Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) { 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_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)) { - 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; @@ -42,6 +39,6 @@ void main() { const float2x3 g = tint_symbol_6(s, 48u); const float3x2 h = tint_symbol_7(s, 80u); 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; } diff --git a/test/buffer/storage/static_index/read.wgsl.expected.msl b/test/buffer/storage/static_index/read.wgsl.expected.msl index d60818e8c5..efc8a38d2d 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.msl +++ b/test/buffer/storage/static_index/read.wgsl.expected.msl @@ -15,12 +15,8 @@ inline vec operator*(packed_vec lhs, matrix rhs) { struct Inner { /* 0x0000 */ int x; }; -struct tint_padded_array_element { - /* 0x0000 */ Inner el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ Inner arr[4]; }; struct S { /* 0x0000 */ packed_int3 a; @@ -33,7 +29,7 @@ struct S { /* 0x0050 */ float3x2 h; /* 0x0068 */ Inner i; /* 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)]]) { diff --git a/test/buffer/storage/static_index/read.wgsl.expected.spvasm b/test/buffer/storage/static_index/read.wgsl.expected.spvasm index 98bf962440..42262ef339 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.spvasm +++ b/test/buffer/storage/static_index/read.wgsl.expected.spvasm @@ -38,7 +38,7 @@ OpMemberDecorate %S 8 Offset 104 OpMemberDecorate %Inner 0 Offset 0 OpMemberDecorate %S 9 Offset 108 - OpDecorate %_arr_Inner_uint_4 ArrayStride 16 + OpDecorate %_arr_Inner_uint_4 ArrayStride 4 OpDecorate %s NonWritable OpDecorate %s Binding 0 OpDecorate %s DescriptorSet 0 diff --git a/test/buffer/storage/static_index/read.wgsl.expected.wgsl b/test/buffer/storage/static_index/read.wgsl.expected.wgsl index 3413a96e30..6c48125089 100644 --- a/test/buffer/storage/static_index/read.wgsl.expected.wgsl +++ b/test/buffer/storage/static_index/read.wgsl.expected.wgsl @@ -12,7 +12,7 @@ struct S { g : mat2x3; h : mat3x2; i : Inner; - j : @stride(16) array; + j : array; } @binding(0) @group(0) var s : S; diff --git a/test/buffer/storage/static_index/write.wgsl b/test/buffer/storage/static_index/write.wgsl index ff18cbc1d1..47827bfc71 100644 --- a/test/buffer/storage/static_index/write.wgsl +++ b/test/buffer/storage/static_index/write.wgsl @@ -12,7 +12,7 @@ struct S { g : mat2x3; h : mat3x2; i : Inner; - j : @stride(16) array; + j : array; }; @binding(0) @group(0) var s : S; @@ -28,5 +28,5 @@ fn main() { s.g = mat2x3(); s.h = mat3x2(); s.i = Inner(); - s.j = @stride(16) array(); + s.j = array(); } diff --git a/test/buffer/storage/static_index/write.wgsl.expected.glsl b/test/buffer/storage/static_index/write.wgsl.expected.glsl index 4434c96ffb..b4db83b30c 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.glsl +++ b/test/buffer/storage/static_index/write.wgsl.expected.glsl @@ -4,9 +4,6 @@ precision mediump float; struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; struct S { ivec3 a; int b; @@ -17,7 +14,7 @@ struct S { mat2x3 g; mat3x2 h; Inner i; - tint_padded_array_element j[4]; + Inner j[4]; }; layout (binding = 0) buffer S_1 { @@ -30,7 +27,7 @@ layout (binding = 0) buffer S_1 { mat2x3 g; mat3x2 h; Inner i; - tint_padded_array_element j[4]; + Inner j[4]; } s; 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); Inner tint_symbol_1 = Inner(0); 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; return; } diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl index 8a2e683a06..f4d174547f 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.hlsl +++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl @@ -1,9 +1,6 @@ struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; 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)); } -void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) { - tint_padded_array_element array[4] = value; +void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, Inner value[4]) { + Inner array[4] = value; { [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)); const Inner tint_symbol_11 = (Inner)0; 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); return; } diff --git a/test/buffer/storage/static_index/write.wgsl.expected.msl b/test/buffer/storage/static_index/write.wgsl.expected.msl index bae3fcd2d2..a3396820ef 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.msl +++ b/test/buffer/storage/static_index/write.wgsl.expected.msl @@ -15,12 +15,8 @@ inline vec operator*(packed_vec lhs, matrix rhs) { struct Inner { /* 0x0000 */ int x; }; -struct tint_padded_array_element { - /* 0x0000 */ Inner el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ Inner arr[4]; }; struct S { /* 0x0000 */ packed_int3 a; @@ -33,7 +29,7 @@ struct S { /* 0x0050 */ float3x2 h; /* 0x0068 */ Inner i; /* 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)]]) { diff --git a/test/buffer/storage/static_index/write.wgsl.expected.spvasm b/test/buffer/storage/static_index/write.wgsl.expected.spvasm index faa44c547b..c099dbda02 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.spvasm +++ b/test/buffer/storage/static_index/write.wgsl.expected.spvasm @@ -38,7 +38,7 @@ OpMemberDecorate %S 8 Offset 104 OpMemberDecorate %Inner 0 Offset 0 OpMemberDecorate %S 9 Offset 108 - OpDecorate %_arr_Inner_uint_4 ArrayStride 16 + OpDecorate %_arr_Inner_uint_4 ArrayStride 4 OpDecorate %s NonReadable OpDecorate %s Binding 0 OpDecorate %s DescriptorSet 0 diff --git a/test/buffer/storage/static_index/write.wgsl.expected.wgsl b/test/buffer/storage/static_index/write.wgsl.expected.wgsl index 595efcecef..68a73e0ad1 100644 --- a/test/buffer/storage/static_index/write.wgsl.expected.wgsl +++ b/test/buffer/storage/static_index/write.wgsl.expected.wgsl @@ -12,7 +12,7 @@ struct S { g : mat2x3; h : mat3x2; i : Inner; - j : @stride(16) array; + j : array; } @binding(0) @group(0) var s : S; @@ -28,5 +28,5 @@ fn main() { s.g = mat2x3(); s.h = mat3x2(); s.i = Inner(); - s.j = @stride(16) array(); + s.j = array(); } diff --git a/test/buffer/uniform/dynamic_index/read.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl index 3e9142baa9..d5b756447a 100644 --- a/test/buffer/uniform/dynamic_index/read.wgsl +++ b/test/buffer/uniform/dynamic_index/read.wgsl @@ -9,7 +9,7 @@ struct Inner { h : vec2; i : mat2x3; @align(16) j : mat3x2; - @align(16) k : @stride(16) array, 4>; + @align(16) k : array, 4>; }; struct S { diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl index 1bee6266de..dbf4c9e6bc 100644 --- a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl +++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl @@ -11,7 +11,7 @@ struct Inner { @align(16) j : mat3x2; @align(16) - k : @stride(16) array, 4>; + k : array, 4>; } struct S { diff --git a/test/buffer/uniform/static_index/read.wgsl b/test/buffer/uniform/static_index/read.wgsl index e1d10e8f2d..6a27ee978e 100644 --- a/test/buffer/uniform/static_index/read.wgsl +++ b/test/buffer/uniform/static_index/read.wgsl @@ -1,5 +1,5 @@ struct Inner { - x : i32; + @size(16) x : i32; }; struct S { @@ -14,7 +14,7 @@ struct S { i : mat2x3; j : mat3x2; @align(16) k : Inner; - @align(16) l : @stride(16) array; + @align(16) l : array; }; @binding(0) @group(0) var s : S; diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.glsl b/test/buffer/uniform/static_index/read.wgsl.expected.glsl index c7b497984e..ee1c6d8d1c 100644 --- a/test/buffer/uniform/static_index/read.wgsl.expected.glsl +++ b/test/buffer/uniform/static_index/read.wgsl.expected.glsl @@ -4,9 +4,6 @@ precision mediump float; struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; struct S { ivec3 a; int b; @@ -19,7 +16,7 @@ struct S { mat2x3 i; mat3x2 j; Inner k; - tint_padded_array_element l[4]; + Inner l[4]; }; layout (binding = 0) uniform S_1 { @@ -34,7 +31,7 @@ layout (binding = 0) uniform S_1 { mat2x3 i; mat3x2 j; Inner k; - tint_padded_array_element l[4]; + Inner l[4]; } s; 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; mat3x2 j = s.j; Inner k = s.k; - tint_padded_array_element l[4] = s.l; + Inner l[4] = s.l; return; } void main() { diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl index 9fa1d0393a..1446321b6b 100644 --- a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl +++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl @@ -1,9 +1,6 @@ struct Inner { int x; }; -struct tint_padded_array_element { - Inner el; -}; cbuffer cbuffer_s : register(b0, space0) { uint4 s[13]; @@ -31,12 +28,12 @@ Inner tint_symbol_10(uint4 buffer[13], uint offset) { 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_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)) { - 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; @@ -55,6 +52,6 @@ void main() { const float2x3 i = tint_symbol_7(s, 64u); const float3x2 j = tint_symbol_8(s, 96u); 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; } diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.msl b/test/buffer/uniform/static_index/read.wgsl.expected.msl index 670ba5d359..46934da129 100644 --- a/test/buffer/uniform/static_index/read.wgsl.expected.msl +++ b/test/buffer/uniform/static_index/read.wgsl.expected.msl @@ -14,13 +14,10 @@ inline vec operator*(packed_vec lhs, matrix rhs) { struct Inner { /* 0x0000 */ int x; -}; -struct tint_padded_array_element { - /* 0x0000 */ Inner el; /* 0x0004 */ int8_t tint_pad[12]; }; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ Inner arr[4]; }; struct S { /* 0x0000 */ packed_int3 a; @@ -35,7 +32,6 @@ struct S { /* 0x0060 */ float3x2 j; /* 0x0078 */ int8_t tint_pad_1[8]; /* 0x0080 */ Inner k; - /* 0x0084 */ int8_t tint_pad_2[12]; /* 0x0090 */ tint_array_wrapper l; }; diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl index b790c19d0c..e23dc3b09f 100644 --- a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl +++ b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl @@ -1,4 +1,5 @@ struct Inner { + @size(16) x : i32; } @@ -16,7 +17,7 @@ struct S { @align(16) k : Inner; @align(16) - l : @stride(16) array; + l : array; } @binding(0) @group(0) var s : S; diff --git a/test/bug/chromium/1273230.wgsl b/test/bug/chromium/1273230.wgsl index 90ed840827..dad0cba28f 100644 --- a/test/bug/chromium/1273230.wgsl +++ b/test/bug/chromium/1273230.wgsl @@ -37,23 +37,23 @@ struct Dbg { }; struct F32s { - values : @stride(4) array; + values : array; }; struct U32s { - values : @stride(4) array; + values : array; }; struct I32s { - values : @stride(4) array; + values : array; }; struct AU32s { - values : @stride(4) array>; + values : array>; }; struct AI32s { - values : @stride(4) array>; + values : array>; }; @binding(0) @group(0) var uniforms : Uniforms; diff --git a/test/bug/chromium/1273230.wgsl.expected.wgsl b/test/bug/chromium/1273230.wgsl.expected.wgsl index 3cdfc2b9da..2d9a2714b2 100644 --- a/test/bug/chromium/1273230.wgsl.expected.wgsl +++ b/test/bug/chromium/1273230.wgsl.expected.wgsl @@ -54,23 +54,23 @@ struct Dbg { } struct F32s { - values : @stride(4) array; + values : array; } struct U32s { - values : @stride(4) array; + values : array; } struct I32s { - values : @stride(4) array; + values : array; } struct AU32s { - values : @stride(4) array>; + values : array>; } struct AI32s { - values : @stride(4) array>; + values : array>; } @binding(0) @group(0) var uniforms : Uniforms; diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl index be2614639a..d05dd5badd 100644 --- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl @@ -1,5 +1,5 @@ struct UBO { - data: @stride(16) array; + data: array, 4>; dynamic_idx: i32; }; @group(0) @binding(0) var ubo: UBO; @@ -10,5 +10,5 @@ struct Result { @stage(compute) @workgroup_size(1) fn f() { - result.out = ubo.data[ubo.dynamic_idx]; + result.out = ubo.data[ubo.dynamic_idx].x; } diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl index cb10e51070..df1ea99b0e 100644 --- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl @@ -1,16 +1,13 @@ #version 310 es precision mediump float; -struct tint_padded_array_element { - int el; -}; struct UBO { - tint_padded_array_element data[4]; + ivec4 data[4]; int dynamic_idx; }; layout (binding = 0) uniform UBO_1 { - tint_padded_array_element data[4]; + ivec4 data[4]; int dynamic_idx; } 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; void f() { - result.tint_symbol = ubo.data[ubo.dynamic_idx].el; + result.tint_symbol = ubo.data[ubo.dynamic_idx].x; return; } void main() { diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl index 3c75f01359..37217139db 100644 --- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl @@ -1,23 +1,20 @@ #include using namespace metal; -struct tint_padded_array_element { - /* 0x0000 */ int el; - /* 0x0004 */ int8_t tint_pad[12]; -}; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ int4 arr[4]; }; struct UBO { /* 0x0000 */ tint_array_wrapper data; /* 0x0040 */ int dynamic_idx; + /* 0x0044 */ int8_t tint_pad[12]; }; struct Result { /* 0x0000 */ int out; }; 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; } diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm index 1026040d3b..8e837ac586 100644 --- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 24 +; Bound: 25 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -17,7 +17,7 @@ OpName %f "f" OpDecorate %UBO Block OpMemberDecorate %UBO 0 Offset 0 - OpDecorate %_arr_int_uint_4 ArrayStride 16 + OpDecorate %_arr_v4int_uint_4 ArrayStride 16 OpMemberDecorate %UBO 1 Offset 64 OpDecorate %ubo NonWritable OpDecorate %ubo DescriptorSet 0 @@ -27,28 +27,29 @@ OpDecorate %result DescriptorSet 0 OpDecorate %result Binding 2 %int = OpTypeInt 32 1 + %v4int = OpTypeVector %int 4 %uint = OpTypeInt 32 0 %uint_4 = OpConstant %uint 4 -%_arr_int_uint_4 = OpTypeArray %int %uint_4 - %UBO = OpTypeStruct %_arr_int_uint_4 %int +%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4 + %UBO = OpTypeStruct %_arr_v4int_uint_4 %int %_ptr_Uniform_UBO = OpTypePointer Uniform %UBO %ubo = OpVariable %_ptr_Uniform_UBO Uniform %Result = OpTypeStruct %int %_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer %void = OpTypeVoid - %11 = OpTypeFunction %void + %12 = OpTypeFunction %void %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %uint_1 = OpConstant %uint 1 %_ptr_Uniform_int = OpTypePointer Uniform %int - %f = OpFunction %void None %11 - %14 = OpLabel - %17 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0 - %20 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1 - %21 = OpLoad %int %20 - %22 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %21 - %23 = OpLoad %int %22 - OpStore %17 %23 + %f = OpFunction %void None %12 + %15 = OpLabel + %18 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0 + %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1 + %22 = OpLoad %int %21 + %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %22 %uint_0 + %24 = OpLoad %int %23 + OpStore %18 %24 OpReturn OpFunctionEnd diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl index 7f3ca0b9b8..0dbf26b277 100644 --- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl +++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl @@ -1,5 +1,5 @@ struct UBO { - data : @stride(16) array; + data : array, 4>; dynamic_idx : i32; } @@ -13,5 +13,5 @@ struct Result { @stage(compute) @workgroup_size(1) fn f() { - result.out = ubo.data[ubo.dynamic_idx]; + result.out = ubo.data[ubo.dynamic_idx].x; } diff --git a/test/bug/tint/1046.wgsl b/test/bug/tint/1046.wgsl index a294edf20d..40eeee6e74 100644 --- a/test/bug/tint/1046.wgsl +++ b/test/bug/tint/1046.wgsl @@ -4,7 +4,7 @@ struct PointLight { }; struct PointLights { - values : @stride(16) array; + values : array; }; struct Uniforms { diff --git a/test/bug/tint/1046.wgsl.expected.wgsl b/test/bug/tint/1046.wgsl.expected.wgsl index b7000e626b..f7ef8d80dd 100644 --- a/test/bug/tint/1046.wgsl.expected.wgsl +++ b/test/bug/tint/1046.wgsl.expected.wgsl @@ -3,7 +3,7 @@ struct PointLight { } struct PointLights { - values : @stride(16) array; + values : array; } struct Uniforms { diff --git a/test/bug/tint/1113.wgsl b/test/bug/tint/1113.wgsl index f3a080a330..87912ffad5 100644 --- a/test/bug/tint/1113.wgsl +++ b/test/bug/tint/1113.wgsl @@ -23,11 +23,11 @@ value_f32_3 : f32; }; - struct F32s { values : @stride(4) array; }; - struct U32s { values : @stride(4) array; }; - struct I32s { values : @stride(4) array; }; - struct AU32s { values : @stride(4) array>; }; - struct AI32s { values : @stride(4) array>; }; + struct F32s { values : array; }; + struct U32s { values : array; }; + struct I32s { values : array; }; + struct AU32s { values : array>; }; + struct AI32s { values : array>; }; // IN @binding(0) @group(0) var uniforms : Uniforms; diff --git a/test/bug/tint/1113.wgsl.expected.wgsl b/test/bug/tint/1113.wgsl.expected.wgsl index 9eeb16c48f..c709577e64 100644 --- a/test/bug/tint/1113.wgsl.expected.wgsl +++ b/test/bug/tint/1113.wgsl.expected.wgsl @@ -23,23 +23,23 @@ struct Dbg { } struct F32s { - values : @stride(4) array; + values : array; } struct U32s { - values : @stride(4) array; + values : array; } struct I32s { - values : @stride(4) array; + values : array; } struct AU32s { - values : @stride(4) array>; + values : array>; } struct AI32s { - values : @stride(4) array>; + values : array>; } @binding(0) @group(0) var uniforms : Uniforms; diff --git a/test/bug/tint/294.wgsl b/test/bug/tint/294.wgsl index 4917c47e70..beff423241 100644 --- a/test/bug/tint/294.wgsl +++ b/test/bug/tint/294.wgsl @@ -3,6 +3,6 @@ struct Light { colour : vec3; }; struct Lights { - light : @stride(32) array; + light : array; }; @group(0) @binding(1) var lights : Lights; diff --git a/test/bug/tint/294.wgsl.expected.wgsl b/test/bug/tint/294.wgsl.expected.wgsl index 87f08def7d..f73ce3ca8d 100644 --- a/test/bug/tint/294.wgsl.expected.wgsl +++ b/test/bug/tint/294.wgsl.expected.wgsl @@ -4,7 +4,7 @@ struct Light { } struct Lights { - light : @stride(32) array; + light : array; } @group(0) @binding(1) var lights : Lights; diff --git a/test/bug/tint/534.wgsl b/test/bug/tint/534.wgsl index 782f193182..77e9f6b12d 100644 --- a/test/bug/tint/534.wgsl +++ b/test/bug/tint/534.wgsl @@ -5,7 +5,7 @@ channelCount : u32; }; struct OutputBuf { - result : @stride(4) array; + result : array; }; @group(0) @binding(0) var src : texture_2d; @group(0) @binding(1) var dst : texture_2d; diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl index dfd0f173b6..7e8620804f 100644 --- a/test/bug/tint/534.wgsl.expected.wgsl +++ b/test/bug/tint/534.wgsl.expected.wgsl @@ -6,7 +6,7 @@ struct Uniforms { } struct OutputBuf { - result : @stride(4) array; + result : array; } @group(0) @binding(0) var src : texture_2d; diff --git a/test/bug/tint/757.wgsl b/test/bug/tint/757.wgsl index b2a05dedaf..24b04184cc 100644 --- a/test/bug/tint/757.wgsl +++ b/test/bug/tint/757.wgsl @@ -7,7 +7,7 @@ @group(0) @binding(1) var myTexture : texture_2d_array; struct Result { - values : @stride(4) array; + values : array; }; @group(0) @binding(3) var result : Result; diff --git a/test/bug/tint/757.wgsl.expected.wgsl b/test/bug/tint/757.wgsl.expected.wgsl index b237cd58cf..6beca1d94d 100644 --- a/test/bug/tint/757.wgsl.expected.wgsl +++ b/test/bug/tint/757.wgsl.expected.wgsl @@ -7,7 +7,7 @@ struct Constants { @group(0) @binding(1) var myTexture : texture_2d_array; struct Result { - values : @stride(4) array; + values : array; } @group(0) @binding(3) var result : Result; diff --git a/test/bug/tint/782.wgsl b/test/bug/tint/782.wgsl index 3fe1f23266..0a4f31b06a 100644 --- a/test/bug/tint/782.wgsl +++ b/test/bug/tint/782.wgsl @@ -1,4 +1,4 @@ -type ArrayExplicitStride = @stride(4) array; +type ArrayExplicitStride = array; type ArrayImplicitStride = array; fn foo() { diff --git a/test/bug/tint/782.wgsl.expected.wgsl b/test/bug/tint/782.wgsl.expected.wgsl index a4f37bcc44..9e895376b1 100644 --- a/test/bug/tint/782.wgsl.expected.wgsl +++ b/test/bug/tint/782.wgsl.expected.wgsl @@ -1,4 +1,4 @@ -type ArrayExplicitStride = @stride(4) array; +type ArrayExplicitStride = array; type ArrayImplicitStride = array; diff --git a/test/bug/tint/922.wgsl b/test/bug/tint/922.wgsl index 450aba182b..1ef80dff43 100644 --- a/test/bug/tint/922.wgsl +++ b/test/bug/tint/922.wgsl @@ -21,12 +21,12 @@ struct ub_SceneParams { }; struct ub_MaterialParams { - u_TexMtx: @stride(32) array; + u_TexMtx: array; u_Misc0_: vec4; }; struct ub_PacketParams { - u_PosMtx: @stride(48) array; + u_PosMtx: array; }; struct VertexOutput { diff --git a/test/bug/tint/922.wgsl.expected.wgsl b/test/bug/tint/922.wgsl.expected.wgsl index ac8b48c438..8151b6e9fb 100644 --- a/test/bug/tint/922.wgsl.expected.wgsl +++ b/test/bug/tint/922.wgsl.expected.wgsl @@ -21,12 +21,12 @@ struct ub_SceneParams { } struct ub_MaterialParams { - u_TexMtx : @stride(32) array; + u_TexMtx : array; u_Misc0_ : vec4; } struct ub_PacketParams { - u_PosMtx : @stride(48) array; + u_PosMtx : array; } struct VertexOutput {