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 <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>

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;
                 },

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<sem::Struct>();
+    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

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
+        // TODO(crbug.com/tint/798) move the @offset attribute handling to the transform::Wgsl
-        // transform::Wgsl sanitizer.
+        // sanitizer.
         if (auto* mem_sem = program_->Sem().Get(mem)) {
             offset = utils::RoundUp(mem_sem->Align(), offset);
             if (uint32_t padding = mem_sem->Offset() - offset) {

test/tint/bug/tint/1088.spvasm.expected.wgsl

@@ -12,8 +12,6 @@ struct LeftOver {
   worldViewProjection : mat4x4<f32>,
   /* @offset(64) */
   time : f32,
-  @size(12)
-  padding : u32,
   /* @offset(80) */
   test2 : Arr,
   /* @offset(208) */

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<f32>,
   /* @offset(48) */

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

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;
+}

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;
+}

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;
+}

test/tint/bug/tint/1776.spvasm.expected.msl

@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+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<int8_t, 12> tint_pad;
+};
+
+struct sb_block {
+  /* 0x0000 */ tint_array<S, 1> 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;
+}
+

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

test/tint/bug/tint/1776.spvasm.expected.wgsl

@@ -0,0 +1,25 @@
+struct S {
+  /* @offset(0) */
+  a : vec4<f32>,
+  /* @offset(16) */
+  b : i32,
+}
+
+type RTArr = array<S>;
+
+struct sb_block {
+  /* @offset(0) */
+  inner : RTArr,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb : sb_block;
+
+fn main_1() {
+  let x_18 : S = sb.inner[1i];
+  return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn main() {
+  main_1();
+}

test/tint/bug/tint/1776.wgsl

@@ -0,0 +1,11 @@
+struct S { /* size: 32 align: 16 */
+  a : vec4<f32>,
+  b : i32,
+}
+
+@group(0) @binding(0) var<storage> sb : array<S>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let x = sb[1];
+}

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;
+}

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;
+}

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;
+}

test/tint/bug/tint/1776.wgsl.expected.msl

@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+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<int8_t, 12> tint_pad;
+};
+
+struct tint_symbol_2 {
+  /* 0x0000 */ tint_array<S, 1> arr;
+};
+
+kernel void tint_symbol(const device tint_symbol_2* tint_symbol_1 [[buffer(0)]]) {
+  S const x = (*(tint_symbol_1)).arr[1];
+  return;
+}
+

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

test/tint/bug/tint/1776.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+struct S {
+  a : vec4<f32>,
+  b : i32,
+}
+
+@group(0) @binding(0) var<storage> sb : array<S>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let x = sb[1];
+}

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 {

test/tint/bug/tint/870.spvasm.expected.msl

@@ -19,6 +19,7 @@ struct sspp962805860buildInformationS {
   /* 0x0010 */ float4 offset;
   /* 0x0020 */ int essence;
   /* 0x0024 */ tint_array<int, 6> orientation;
+  /* 0x003c */ tint_array<int8_t, 4> tint_pad;
 };
 
 struct x_B4_BuildInformation {