Padding tests for vec3h in arrays (#3985)

Adds tests that the padding between vec3h elements in an array is not
disturbed.

src/webgpu/listing_meta.json

@@ -1863,6 +1863,8 @@
   "webgpu:shader,execution,padding:array_of_matCx3:*": { "subcaseMS": 8.650 },
   "webgpu:shader,execution,padding:array_of_struct:*": { "subcaseMS": 5.801 },
   "webgpu:shader,execution,padding:array_of_vec3:*": { "subcaseMS": 10.500 },
+  "webgpu:shader,execution,padding:array_of_vec3h,elementwise:*": { "subcaseMS": 24.607 },
+  "webgpu:shader,execution,padding:array_of_vec3h:*": { "subcaseMS": 26.941 },
   "webgpu:shader,execution,padding:matCx3:*": { "subcaseMS": 10.050 },
   "webgpu:shader,execution,padding:struct_explicit:*": { "subcaseMS": 12.000 },
   "webgpu:shader,execution,padding:struct_implicit:*": { "subcaseMS": 33.201 },

src/webgpu/shader/execution/padding.spec.ts

@@ -263,6 +263,87 @@ g.test('array_of_vec3')
     );
   });
 
+g.test('array_of_vec3h')
+  .desc(
+    `Test that padding bytes in between array elements are preserved when f16 elements are used.
+
+     This test defines creates a read-write storage buffer with type array<vec3h, 4>. The shader
+     assigns the whole variable at once, and we then test that data in the padding bytes was
+     preserved.
+    `
+  )
+  .beforeAllSubcases(t => {
+    t.selectDeviceOrSkipTestCase('shader-f16');
+  })
+  .fn(t => {
+    const wgsl = `
+      enable f16;
+      @group(0) @binding(0) var<storage, read_write> buffer : array<vec3<f16>, 4>;
+
+      @compute @workgroup_size(1)
+      fn main() {
+        buffer = array<vec3<f16>, 4>(
+          vec3(1h),
+          vec3(2h),
+          vec3(3h),
+          vec3(4h),
+        );
+      }
+    `;
+    runShaderTest(
+      t,
+      wgsl,
+      new Uint32Array([
+        // buffer[0]
+        0x3c003c00, 0xdead3c00,
+        // buffer[1]
+        0x40004000, 0xdead4000,
+        // buffer[2]
+        0x42004200, 0xdead4200,
+        // buffer[2]
+        0x44004400, 0xdead4400,
+      ])
+    );
+  });
+
+g.test('array_of_vec3h,elementwise')
+  .desc(
+    `Test that padding bytes in between array elements are preserved when f16 elements are used.
+
+     This test defines creates a read-write storage buffer with type array<vec3h, 4>. The shader
+     assigns one element per thread, and we then test that data in the padding bytes was
+     preserved.
+    `
+  )
+  .beforeAllSubcases(t => {
+    t.selectDeviceOrSkipTestCase('shader-f16');
+  })
+  .fn(t => {
+    const wgsl = `
+      enable f16;
+      @group(0) @binding(0) var<storage, read_write> buffer : array<vec3<f16>>;
+
+      @compute @workgroup_size(4)
+      fn main(@builtin(local_invocation_index) lid : u32) {
+        buffer[lid] = vec3h(f16(lid + 1));
+      }
+    `;
+    runShaderTest(
+      t,
+      wgsl,
+      new Uint32Array([
+        // buffer[0]
+        0x3c003c00, 0xdead3c00,
+        // buffer[1]
+        0x40004000, 0xdead4000,
+        // buffer[2]
+        0x42004200, 0xdead4200,
+        // buffer[2]
+        0x44004400, 0xdead4400,
+      ])
+    );
+  });
+
 g.test('array_of_struct')
   .desc(
     `Test that padding bytes in between array elements are preserved.