From d257e287929813fae395a2d6766fe305d883ce1e Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 1 Dec 2022 18:41:57 +0000 Subject: [PATCH] tint: Don't override alignment with @offset Overriding the alignment to 1 would cause nested structures to be incorrectly laid out. The fix: Don't override the alignment. All struct layout validation works on the sem offsets, so none of this has to change. Bug: tint:1776 Change-Id: Ic01d45fb2790cd823ed9a55e336860ebdc351aea Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/112603 Kokoro: Ben Clayton Reviewed-by: Dan Sinclair Commit-Queue: Ben Clayton --- src/tint/resolver/resolver.cc | 1 - src/tint/resolver/struct_layout_test.cc | 42 ++++++++++++++++ src/tint/writer/wgsl/generator_impl.cc | 4 +- test/tint/bug/tint/1088.spvasm.expected.wgsl | 2 - test/tint/bug/tint/1520.spvasm.expected.wgsl | 2 - test/tint/bug/tint/1776.spvasm | 44 +++++++++++++++++ .../bug/tint/1776.spvasm.expected.dxc.hlsl | 22 +++++++++ .../bug/tint/1776.spvasm.expected.fxc.hlsl | 22 +++++++++ test/tint/bug/tint/1776.spvasm.expected.glsl | 28 +++++++++++ test/tint/bug/tint/1776.spvasm.expected.msl | 36 ++++++++++++++ .../tint/bug/tint/1776.spvasm.expected.spvasm | 49 +++++++++++++++++++ test/tint/bug/tint/1776.spvasm.expected.wgsl | 25 ++++++++++ test/tint/bug/tint/1776.wgsl | 11 +++++ .../tint/bug/tint/1776.wgsl.expected.dxc.hlsl | 17 +++++++ .../tint/bug/tint/1776.wgsl.expected.fxc.hlsl | 17 +++++++ test/tint/bug/tint/1776.wgsl.expected.glsl | 23 +++++++++ test/tint/bug/tint/1776.wgsl.expected.msl | 31 ++++++++++++ test/tint/bug/tint/1776.wgsl.expected.spvasm | 44 +++++++++++++++++ test/tint/bug/tint/1776.wgsl.expected.wgsl | 11 +++++ test/tint/bug/tint/870.spvasm.expected.glsl | 1 + test/tint/bug/tint/870.spvasm.expected.msl | 1 + 21 files changed, 426 insertions(+), 7 deletions(-) create mode 100644 test/tint/bug/tint/1776.spvasm create mode 100644 test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1776.spvasm.expected.glsl create mode 100644 test/tint/bug/tint/1776.spvasm.expected.msl create mode 100644 test/tint/bug/tint/1776.spvasm.expected.spvasm create mode 100644 test/tint/bug/tint/1776.spvasm.expected.wgsl create mode 100644 test/tint/bug/tint/1776.wgsl create mode 100644 test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1776.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1776.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1776.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1776.wgsl.expected.wgsl diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc index 3ad86617fd..60f8207798 100644 --- a/src/tint/resolver/resolver.cc +++ b/src/tint/resolver/resolver.cc @@ -3173,7 +3173,6 @@ sem::Struct* Resolver::Structure(const ast::Struct* str) { AddError("offsets must be in ascending order", o->source); return false; } - align = 1; has_offset_attr = true; return true; }, diff --git a/src/tint/resolver/struct_layout_test.cc b/src/tint/resolver/struct_layout_test.cc index 1e5ec45464..acfe6916c3 100644 --- a/src/tint/resolver/struct_layout_test.cc +++ b/src/tint/resolver/struct_layout_test.cc @@ -555,5 +555,47 @@ TEST_F(ResolverStructLayoutTest, StructWithLotsOfPadding) { } } +TEST_F(ResolverStructLayoutTest, OffsetAttributes) { + auto* inner = Structure("Inner", utils::Vector{ + Member("a", ty.f32(), utils::Vector{MemberOffset(8_i)}), + Member("b", ty.f32(), utils::Vector{MemberOffset(16_i)}), + Member("c", ty.f32(), utils::Vector{MemberOffset(32_i)}), + }); + auto* s = Structure("S", utils::Vector{ + Member("a", ty.f32(), utils::Vector{MemberOffset(4_i)}), + Member("b", ty.u32(), utils::Vector{MemberOffset(8_i)}), + Member("c", ty.Of(inner), utils::Vector{MemberOffset(32_i)}), + Member("d", ty.i32()), + Member("e", ty.i32(), utils::Vector{MemberOffset(128_i)}), + }); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + + auto* sem = TypeOf(s)->As(); + ASSERT_NE(sem, nullptr); + EXPECT_EQ(sem->Size(), 132u); + EXPECT_EQ(sem->SizeNoPadding(), 132u); + EXPECT_EQ(sem->Align(), 4u); + ASSERT_EQ(sem->Members().size(), 5u); + EXPECT_EQ(sem->Members()[0]->Offset(), 4u); + EXPECT_EQ(sem->Members()[0]->Align(), 4u); + EXPECT_EQ(sem->Members()[0]->Size(), 4u); + EXPECT_EQ(sem->Members()[1]->Offset(), 8u); + EXPECT_EQ(sem->Members()[1]->Align(), 4u); + EXPECT_EQ(sem->Members()[1]->Size(), 4u); + EXPECT_EQ(sem->Members()[2]->Offset(), 32u); + EXPECT_EQ(sem->Members()[2]->Align(), 4u); + EXPECT_EQ(sem->Members()[2]->Size(), 36u); + EXPECT_EQ(sem->Members()[3]->Offset(), 68u); + EXPECT_EQ(sem->Members()[3]->Align(), 4u); + EXPECT_EQ(sem->Members()[3]->Size(), 4u); + EXPECT_EQ(sem->Members()[4]->Offset(), 128u); + EXPECT_EQ(sem->Members()[4]->Align(), 4u); + EXPECT_EQ(sem->Members()[4]->Size(), 4u); + for (auto& m : sem->Members()) { + EXPECT_EQ(m->Struct()->Declaration(), s); + } +} + } // namespace } // namespace tint::resolver diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc index f701146347..cfd512b1f7 100644 --- a/src/tint/writer/wgsl/generator_impl.cc +++ b/src/tint/writer/wgsl/generator_impl.cc @@ -606,8 +606,8 @@ bool GeneratorImpl::EmitStructType(const ast::Struct* str) { increment_indent(); uint32_t offset = 0; for (auto* mem : str->members) { - // TODO(crbug.com/tint/798) move the @offset attribute handling to the - // transform::Wgsl sanitizer. + // TODO(crbug.com/tint/798) move the @offset attribute handling to the transform::Wgsl + // sanitizer. if (auto* mem_sem = program_->Sem().Get(mem)) { offset = utils::RoundUp(mem_sem->Align(), offset); if (uint32_t padding = mem_sem->Offset() - offset) { diff --git a/test/tint/bug/tint/1088.spvasm.expected.wgsl b/test/tint/bug/tint/1088.spvasm.expected.wgsl index 54ce4f293c..54683995b5 100644 --- a/test/tint/bug/tint/1088.spvasm.expected.wgsl +++ b/test/tint/bug/tint/1088.spvasm.expected.wgsl @@ -12,8 +12,6 @@ struct LeftOver { worldViewProjection : mat4x4, /* @offset(64) */ time : f32, - @size(12) - padding : u32, /* @offset(80) */ test2 : Arr, /* @offset(208) */ diff --git a/test/tint/bug/tint/1520.spvasm.expected.wgsl b/test/tint/bug/tint/1520.spvasm.expected.wgsl index 931944d3e6..812b115b00 100644 --- a/test/tint/bug/tint/1520.spvasm.expected.wgsl +++ b/test/tint/bug/tint/1520.spvasm.expected.wgsl @@ -3,8 +3,6 @@ struct UniformBuffer { padding : u32, /* @offset(16) */ unknownInput_S1_c0 : f32, - @size(12) - padding_1 : u32, /* @offset(32) */ ucolorRed_S1_c0 : vec4, /* @offset(48) */ diff --git a/test/tint/bug/tint/1776.spvasm b/test/tint/bug/tint/1776.spvasm new file mode 100644 index 0000000000..3e85c9ed17 --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 19 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %sb_block "sb_block" + OpMemberName %sb_block 0 "inner" + OpName %S "S" + OpMemberName %S 0 "a" + OpMemberName %S 1 "b" + OpName %sb "sb" + OpName %main "main" + OpDecorate %sb_block Block + OpMemberDecorate %sb_block 0 Offset 0 + OpMemberDecorate %S 0 Offset 0 + OpMemberDecorate %S 1 Offset 16 + OpDecorate %_runtimearr_S ArrayStride 32 + OpDecorate %sb NonWritable + OpDecorate %sb DescriptorSet 0 + OpDecorate %sb Binding 0 + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 + %int = OpTypeInt 32 1 + %S = OpTypeStruct %v4float %int +%_runtimearr_S = OpTypeRuntimeArray %S + %sb_block = OpTypeStruct %_runtimearr_S +%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block + %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %int_1 = OpConstant %int 1 +%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %main = OpFunction %void None %9 + %12 = OpLabel + %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1 + %18 = OpLoad %S %17 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl new file mode 100644 index 0000000000..1cd2986edb --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl @@ -0,0 +1,22 @@ +struct S { + float4 a; + int b; +}; + +RWByteAddressBuffer sb : register(u0, space0); + +S tint_symbol(RWByteAddressBuffer buffer, uint offset) { + const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))}; + return tint_symbol_3; +} + +void main_1() { + const S x_18 = tint_symbol(sb, 32u); + return; +} + +[numthreads(1, 1, 1)] +void main() { + main_1(); + return; +} diff --git a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl new file mode 100644 index 0000000000..1cd2986edb --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl @@ -0,0 +1,22 @@ +struct S { + float4 a; + int b; +}; + +RWByteAddressBuffer sb : register(u0, space0); + +S tint_symbol(RWByteAddressBuffer buffer, uint offset) { + const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))}; + return tint_symbol_3; +} + +void main_1() { + const S x_18 = tint_symbol(sb, 32u); + return; +} + +[numthreads(1, 1, 1)] +void main() { + main_1(); + return; +} diff --git a/test/tint/bug/tint/1776.spvasm.expected.glsl b/test/tint/bug/tint/1776.spvasm.expected.glsl new file mode 100644 index 0000000000..e73e547ed0 --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.glsl @@ -0,0 +1,28 @@ +#version 310 es + +struct S { + vec4 a; + int b; + uint pad; + uint pad_1; + uint pad_2; +}; + +layout(binding = 0, std430) buffer sb_block_ssbo { + S inner[]; +} sb; + +void main_1() { + S x_18 = sb.inner[1]; + return; +} + +void tint_symbol() { + main_1(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1776.spvasm.expected.msl b/test/tint/bug/tint/1776.spvasm.expected.msl new file mode 100644 index 0000000000..0933ac03b3 --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.msl @@ -0,0 +1,36 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +struct S { + /* 0x0000 */ float4 a; + /* 0x0010 */ int b; + /* 0x0014 */ tint_array tint_pad; +}; + +struct sb_block { + /* 0x0000 */ tint_array inner; +}; + +void main_1(device sb_block* const tint_symbol_1) { + S const x_18 = (*(tint_symbol_1)).inner[1]; + return; +} + +kernel void tint_symbol(device sb_block* tint_symbol_2 [[buffer(0)]]) { + main_1(tint_symbol_2); + return; +} + diff --git a/test/tint/bug/tint/1776.spvasm.expected.spvasm b/test/tint/bug/tint/1776.spvasm.expected.spvasm new file mode 100644 index 0000000000..6e076d1860 --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.spvasm @@ -0,0 +1,49 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %sb_block "sb_block" + OpMemberName %sb_block 0 "inner" + OpName %S "S" + OpMemberName %S 0 "a" + OpMemberName %S 1 "b" + OpName %sb "sb" + OpName %main_1 "main_1" + OpName %main "main" + OpDecorate %sb_block Block + OpMemberDecorate %sb_block 0 Offset 0 + OpMemberDecorate %S 0 Offset 0 + OpMemberDecorate %S 1 Offset 16 + OpDecorate %_runtimearr_S ArrayStride 32 + OpDecorate %sb DescriptorSet 0 + OpDecorate %sb Binding 0 + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 + %int = OpTypeInt 32 1 + %S = OpTypeStruct %v4float %int +%_runtimearr_S = OpTypeRuntimeArray %S + %sb_block = OpTypeStruct %_runtimearr_S +%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block + %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %int_1 = OpConstant %int 1 +%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %main_1 = OpFunction %void None %9 + %12 = OpLabel + %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1 + %18 = OpLoad %S %17 + OpReturn + OpFunctionEnd + %main = OpFunction %void None %9 + %20 = OpLabel + %21 = OpFunctionCall %void %main_1 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1776.spvasm.expected.wgsl b/test/tint/bug/tint/1776.spvasm.expected.wgsl new file mode 100644 index 0000000000..c5f6bc7264 --- /dev/null +++ b/test/tint/bug/tint/1776.spvasm.expected.wgsl @@ -0,0 +1,25 @@ +struct S { + /* @offset(0) */ + a : vec4, + /* @offset(16) */ + b : i32, +} + +type RTArr = array; + +struct sb_block { + /* @offset(0) */ + inner : RTArr, +} + +@group(0) @binding(0) var sb : sb_block; + +fn main_1() { + let x_18 : S = sb.inner[1i]; + return; +} + +@compute @workgroup_size(1i, 1i, 1i) +fn main() { + main_1(); +} diff --git a/test/tint/bug/tint/1776.wgsl b/test/tint/bug/tint/1776.wgsl new file mode 100644 index 0000000000..893b64ca87 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl @@ -0,0 +1,11 @@ +struct S { /* size: 32 align: 16 */ + a : vec4, + b : i32, +} + +@group(0) @binding(0) var sb : array; + +@compute @workgroup_size(1) +fn main() { + let x = sb[1]; +} diff --git a/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..36f87a7804 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl @@ -0,0 +1,17 @@ +struct S { + float4 a; + int b; +}; + +ByteAddressBuffer sb : register(t0, space0); + +S tint_symbol(ByteAddressBuffer buffer, uint offset) { + const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))}; + return tint_symbol_3; +} + +[numthreads(1, 1, 1)] +void main() { + const S x = tint_symbol(sb, 32u); + return; +} diff --git a/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..36f87a7804 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl @@ -0,0 +1,17 @@ +struct S { + float4 a; + int b; +}; + +ByteAddressBuffer sb : register(t0, space0); + +S tint_symbol(ByteAddressBuffer buffer, uint offset) { + const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))}; + return tint_symbol_3; +} + +[numthreads(1, 1, 1)] +void main() { + const S x = tint_symbol(sb, 32u); + return; +} diff --git a/test/tint/bug/tint/1776.wgsl.expected.glsl b/test/tint/bug/tint/1776.wgsl.expected.glsl new file mode 100644 index 0000000000..6a83ec28e8 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.glsl @@ -0,0 +1,23 @@ +#version 310 es + +struct S { + vec4 a; + int b; + uint pad; + uint pad_1; + uint pad_2; +}; + +layout(binding = 0, std430) buffer sb_block_ssbo { + S inner[]; +} sb; + +void tint_symbol() { + S x = sb.inner[1]; +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1776.wgsl.expected.msl b/test/tint/bug/tint/1776.wgsl.expected.msl new file mode 100644 index 0000000000..38097d0784 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.msl @@ -0,0 +1,31 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +struct S { + /* 0x0000 */ float4 a; + /* 0x0010 */ int b; + /* 0x0014 */ tint_array tint_pad; +}; + +struct tint_symbol_2 { + /* 0x0000 */ tint_array arr; +}; + +kernel void tint_symbol(const device tint_symbol_2* tint_symbol_1 [[buffer(0)]]) { + S const x = (*(tint_symbol_1)).arr[1]; + return; +} + diff --git a/test/tint/bug/tint/1776.wgsl.expected.spvasm b/test/tint/bug/tint/1776.wgsl.expected.spvasm new file mode 100644 index 0000000000..3e85c9ed17 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.spvasm @@ -0,0 +1,44 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 19 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %sb_block "sb_block" + OpMemberName %sb_block 0 "inner" + OpName %S "S" + OpMemberName %S 0 "a" + OpMemberName %S 1 "b" + OpName %sb "sb" + OpName %main "main" + OpDecorate %sb_block Block + OpMemberDecorate %sb_block 0 Offset 0 + OpMemberDecorate %S 0 Offset 0 + OpMemberDecorate %S 1 Offset 16 + OpDecorate %_runtimearr_S ArrayStride 32 + OpDecorate %sb NonWritable + OpDecorate %sb DescriptorSet 0 + OpDecorate %sb Binding 0 + %float = OpTypeFloat 32 + %v4float = OpTypeVector %float 4 + %int = OpTypeInt 32 1 + %S = OpTypeStruct %v4float %int +%_runtimearr_S = OpTypeRuntimeArray %S + %sb_block = OpTypeStruct %_runtimearr_S +%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block + %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer + %void = OpTypeVoid + %9 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %int_1 = OpConstant %int 1 +%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %main = OpFunction %void None %9 + %12 = OpLabel + %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1 + %18 = OpLoad %S %17 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1776.wgsl.expected.wgsl b/test/tint/bug/tint/1776.wgsl.expected.wgsl new file mode 100644 index 0000000000..08733dcc24 --- /dev/null +++ b/test/tint/bug/tint/1776.wgsl.expected.wgsl @@ -0,0 +1,11 @@ +struct S { + a : vec4, + b : i32, +} + +@group(0) @binding(0) var sb : array; + +@compute @workgroup_size(1) +fn main() { + let x = sb[1]; +} diff --git a/test/tint/bug/tint/870.spvasm.expected.glsl b/test/tint/bug/tint/870.spvasm.expected.glsl index 0cddbf5f3b..f56b3debd8 100644 --- a/test/tint/bug/tint/870.spvasm.expected.glsl +++ b/test/tint/bug/tint/870.spvasm.expected.glsl @@ -6,6 +6,7 @@ struct sspp962805860buildInformationS { vec4 offset; int essence; int orientation[6]; + uint pad; }; struct x_B4_BuildInformation { diff --git a/test/tint/bug/tint/870.spvasm.expected.msl b/test/tint/bug/tint/870.spvasm.expected.msl index 9a97441efd..408a97122c 100644 --- a/test/tint/bug/tint/870.spvasm.expected.msl +++ b/test/tint/bug/tint/870.spvasm.expected.msl @@ -19,6 +19,7 @@ struct sspp962805860buildInformationS { /* 0x0010 */ float4 offset; /* 0x0020 */ int essence; /* 0x0024 */ tint_array orientation; + /* 0x003c */ tint_array tint_pad; }; struct x_B4_BuildInformation {