tint/msl: Preserve trailing vec3 padding

In order to preserve padding properly for MSL, we need to use its
packed_vec type for all vec3 types in storage buffers, not just struct
members. This commit includes a complete rewrite of the PackedVec3
transform to achieve this. The key details are:

* An internal `__packed_vec3<>` type was added, which corresponds to a
  `type::Vector` with an additional flag to indicate that it will be
  emitted as packed vector.

* The `PackedVec3` transform replaces all vec3 types used in
  host-shareable address spaces with the internal `__packed_vec3`
  type. This includes vec3 types that appear as the store type of a
  pointer.

* When used as an array element, these `__packed_vec3` types are
  wrapped in a struct that contains a single `__packed_vec3`
  member. This allows us to add an `@align()` attribute that ensures
  that `array<vec3<T>>` still has the correct array element stride.

* When the `vec3<T>` appears as a struct member in the input program,
  we apply the `@align()` to that member to ensure that we do not
  change its offset.

* Matrix types with three rows that are used in memory are replaced
  with an array of columns, where each column uses a `__packed_vec3`
  inside an aligned wrapper structure as above.

* Accesses to host-shareable memory that involve any of these types
  invoke a "pack" or "unpack" helper function to convert them to the
  equivalent type that uses `__packed_vec3` or a regular `vec3` as
  required.

* The `chromium_internal_relaxed_uniform_layout` extension is used to
  avoid issues where modifying a type in the uniform address space
  triggers stricter layout validation rules.

Bug: tint:1571
Fixed: tint:1837
Change-Id: Idaf2da2f5bcb2be00c85ec657edfb614186476bb
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/121200
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.msl

@@ -1,17 +1,43 @@
 #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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
 int i() {
   thread int tint_symbol_1 = 0;
   tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
   return tint_symbol_1;
 }
 
-kernel void f(const constant half4x3* tint_symbol_2 [[buffer(0)]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_2 [[buffer(0)]]) {
   int const tint_symbol = i();
   int const p_m_i_save = tint_symbol;
-  half4x3 const l_m = *(tint_symbol_2);
-  half3 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
+  half4x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_2));
+  half3 const l_m_i = half3((*(tint_symbol_2))[p_m_i_save].elements);
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/static_index_via_ptr.wgsl.expected.msl

@@ -1,15 +1,41 @@
 #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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
 int i() {
   thread int tint_symbol = 0;
   tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
   return tint_symbol;
 }
 
-kernel void f(const constant half4x3* tint_symbol_1 [[buffer(0)]]) {
-  half4x3 const l_m = *(tint_symbol_1);
-  half3 const l_m_1 = (*(tint_symbol_1))[1];
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_1 [[buffer(0)]]) {
+  half4x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
+  half3 const l_m_1 = half3((*(tint_symbol_1))[1].elements);
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_builtin.wgsl.expected.msl

@@ -1,10 +1,36 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void f(const constant half4x3* tint_symbol [[buffer(0)]]) {
-  half3x4 const t = transpose(*(tint_symbol));
-  half const l = length((*(tint_symbol))[1]);
-  half const a = fabs(half3((*(tint_symbol))[0]).zxy[0]);
+
+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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol [[buffer(0)]]) {
+  half3x4 const t = transpose(tint_unpack_vec3_in_composite(*(tint_symbol)));
+  half const l = length(half3((*(tint_symbol))[1].elements));
+  half const a = fabs(half3((*(tint_symbol))[0].elements).zxy[0]);
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_fn.wgsl.expected.msl

@@ -1,6 +1,32 @@
 #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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
 void a(half4x3 m) {
 }
 
@@ -10,12 +36,12 @@ void b(half3 v) {
 void c(half f_1) {
 }
 
-kernel void f(const constant half4x3* tint_symbol [[buffer(0)]]) {
-  a(*(tint_symbol));
-  b((*(tint_symbol))[1]);
-  b(half3((*(tint_symbol))[1]).zxy);
-  c((*(tint_symbol))[1][0]);
-  c(half3((*(tint_symbol))[1]).zxy[0]);
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol [[buffer(0)]]) {
+  a(tint_unpack_vec3_in_composite(*(tint_symbol)));
+  b(half3((*(tint_symbol))[1].elements));
+  b(half3((*(tint_symbol))[1].elements).zxy);
+  c((*(tint_symbol))[1].elements[0]);
+  c(half3((*(tint_symbol))[1].elements).zxy[0]);
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_private.wgsl.expected.msl

@@ -1,12 +1,38 @@
 #include <metal_stdlib>
 
 using namespace metal;
-kernel void f(const constant half4x3* tint_symbol_1 [[buffer(0)]]) {
+
+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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_1 [[buffer(0)]]) {
   thread half4x3 tint_symbol = half4x3(0.0h);
-  tint_symbol = *(tint_symbol_1);
-  tint_symbol[1] = (*(tint_symbol_1))[0];
-  tint_symbol[1] = half3((*(tint_symbol_1))[0]).zxy;
-  tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
+  tint_symbol = tint_unpack_vec3_in_composite(*(tint_symbol_1));
+  tint_symbol[1] = half3((*(tint_symbol_1))[0].elements);
+  tint_symbol[1] = half3((*(tint_symbol_1))[0].elements).zxy;
+  tint_symbol[0][1] = (*(tint_symbol_1))[1].elements[0];
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_storage.wgsl.expected.msl

@@ -1,18 +1,44 @@
 #include <metal_stdlib>
 
 using namespace metal;
-void assign_and_preserve_padding(device half4x3* const dest, half4x3 value) {
-  (*(dest))[0u] = value[0u];
-  (*(dest))[1u] = value[1u];
-  (*(dest))[2u] = value[2u];
-  (*(dest))[3u] = value[3u];
+
+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 tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
 }
 
-kernel void f(device half4x3* tint_symbol [[buffer(1)]], const constant half4x3* tint_symbol_1 [[buffer(0)]]) {
-  assign_and_preserve_padding(tint_symbol, *(tint_symbol_1));
-  (*(tint_symbol))[1] = (*(tint_symbol_1))[0];
-  (*(tint_symbol))[1] = half3((*(tint_symbol_1))[0]).zxy;
-  (*(tint_symbol))[0][1] = (*(tint_symbol_1))[1][0];
+void assign_and_preserve_padding(device tint_array<tint_packed_vec3_f16_array_element, 4>* const dest, half4x3 value) {
+  (*(dest))[0u].elements = packed_half3(value[0u]);
+  (*(dest))[1u].elements = packed_half3(value[1u]);
+  (*(dest))[2u].elements = packed_half3(value[2u]);
+  (*(dest))[3u].elements = packed_half3(value[3u]);
+}
+
+kernel void f(device tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol [[buffer(1)]], const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_1 [[buffer(0)]]) {
+  assign_and_preserve_padding(tint_symbol, tint_unpack_vec3_in_composite(*(tint_symbol_1)));
+  (*(tint_symbol))[1].elements = (*(tint_symbol_1))[0].elements;
+  (*(tint_symbol))[1].elements = packed_half3(half3((*(tint_symbol_1))[0].elements).zxy);
+  (*(tint_symbol))[0].elements[1] = (*(tint_symbol_1))[1].elements[0];
   return;
 }
 

test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.msl

@@ -1,22 +1,48 @@
 #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 tint_symbol_5 {
   half4x3 w;
 };
 
-void f_inner(uint local_invocation_index, threadgroup half4x3* const tint_symbol, const constant half4x3* const tint_symbol_1) {
+struct tint_packed_vec3_f16_array_element {
+  /* 0x0000 */ packed_half3 elements;
+  /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
+};
+
+half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_element, 4> in) {
+  half4x3 result = half4x3(0.0h);
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    result[i] = half3(in[i].elements);
+  }
+  return result;
+}
+
+void f_inner(uint local_invocation_index, threadgroup half4x3* const tint_symbol, const constant tint_array<tint_packed_vec3_f16_array_element, 4>* const tint_symbol_1) {
   {
     *(tint_symbol) = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  *(tint_symbol) = *(tint_symbol_1);
-  (*(tint_symbol))[1] = (*(tint_symbol_1))[0];
-  (*(tint_symbol))[1] = half3((*(tint_symbol_1))[0]).zxy;
-  (*(tint_symbol))[0][1] = (*(tint_symbol_1))[1][0];
+  *(tint_symbol) = tint_unpack_vec3_in_composite(*(tint_symbol_1));
+  (*(tint_symbol))[1] = half3((*(tint_symbol_1))[0].elements);
+  (*(tint_symbol))[1] = half3((*(tint_symbol_1))[0].elements).zxy;
+  (*(tint_symbol))[0][1] = (*(tint_symbol_1))[1].elements[0];
 }
 
-kernel void f(const constant half4x3* tint_symbol_4 [[buffer(0)]], threadgroup tint_symbol_5* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 4>* tint_symbol_4 [[buffer(0)]], threadgroup tint_symbol_5* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
   threadgroup half4x3* const tint_symbol_2 = &((*(tint_symbol_3)).w);
   f_inner(local_invocation_index, tint_symbol_2, tint_symbol_4);
   return;