From ae641a6d381518745d0161c631227b86f33d7ebb Mon Sep 17 00:00:00 2001
From: alan-baker <alanbaker@google.com>
Date: Mon, 7 Oct 2024 10:40:13 -0400
Subject: [PATCH] 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                |  2 +
 src/webgpu/shader/execution/padding.spec.ts | 81 +++++++++++++++++++++
 2 files changed, 83 insertions(+)

diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index 86ed08912622..c7d10a4c0f26 100644
--- a/src/webgpu/listing_meta.json
+++ b/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 },
diff --git a/src/webgpu/shader/execution/padding.spec.ts b/src/webgpu/shader/execution/padding.spec.ts
index 3a3671bcc3ff..c9e230013590 100644
--- a/src/webgpu/shader/execution/padding.spec.ts
+++ b/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.