From d4a925c48832d7d09ecde02246e14156f54a4c17 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Fran=C3=A7ois=20Beaufort?= <beaufort.francois@gmail.com>
Date: Mon, 29 Jan 2024 15:41:56 +0100
Subject: [PATCH] Add shared worker support

---
 docs/intro/developing.md                      |   5 +-
 src/common/internal/query/query.ts            |   6 +-
 src/common/runtime/helper/options.ts          |  14 +-
 .../runtime/helper/test_worker-worker.ts      |  18 +-
 src/common/runtime/helper/test_worker.ts      |  49 +++-
 src/common/runtime/standalone.ts              |  12 +-
 src/common/runtime/wpt.ts                     |  15 +-
 src/common/tools/gen_wpt_cts_html.ts          |  12 +-
 src/resources/cache/hashes.json               | 214 +++++++++---------
 .../cache/webgpu/shader/execution/bitcast.bin | Bin 2221448 -> 2221448 bytes
 src/webgpu/listing_meta.json                  | 159 ++++++-------
 src/webgpu/web_platform/worker/worker.spec.ts |  34 ++-
 src/webgpu/web_platform/worker/worker.ts      |  20 +-
 .../web_platform/worker/worker_launcher.ts    |  19 +-
 14 files changed, 351 insertions(+), 226 deletions(-)

diff --git a/docs/intro/developing.md b/docs/intro/developing.md
index a942a6842d75..d8869f985b31 100644
--- a/docs/intro/developing.md
+++ b/docs/intro/developing.md
@@ -48,7 +48,7 @@ You can use this to preview how your test plan will appear.
 You can view different suites (webgpu, unittests, stress, etc.) or different subtrees of
 the test suite.
 
-- `http://localhost:8080/standalone/` (defaults to `?runnow=0&worker=0&debug=0&q=webgpu:*`)
+- `http://localhost:8080/standalone/` (defaults to `?runnow=0&debug=0&q=webgpu:*`)
 - `http://localhost:8080/standalone/?q=unittests:*`
 - `http://localhost:8080/standalone/?q=unittests:basic:*`
 
@@ -56,7 +56,8 @@ The following url parameters change how the harness runs:
 
 - `runnow=1` runs all matching tests on page load.
 - `debug=1` enables verbose debug logging from tests.
-- `worker=1` runs the tests on a Web Worker instead of the main thread.
+- `worker=dedicated` runs the tests on a dedicated worker instead of the main thread.
+- `worker=shared` runs the tests on a shared worker instead of the main thread.
 - `power_preference=low-power` runs most tests passing `powerPreference: low-power` to `requestAdapter`
 - `power_preference=high-performance` runs most tests passing `powerPreference: high-performance` to `requestAdapter`
 
diff --git a/src/common/internal/query/query.ts b/src/common/internal/query/query.ts
index 7c72a62f885a..ad7cf246e9d9 100644
--- a/src/common/internal/query/query.ts
+++ b/src/common/internal/query/query.ts
@@ -1,5 +1,5 @@
 import { TestParams } from '../../framework/fixture.js';
-import { optionEnabled } from '../../runtime/helper/options.js';
+import { optionString } from '../../runtime/helper/options.js';
 import { assert, unreachable } from '../../util/util.js';
 import { Expectation } from '../logging/result.js';
 
@@ -188,12 +188,12 @@ export function parseExpectationsForTestQuery(
       assert(
         expectationURL.pathname === wptURL.pathname,
         `Invalid expectation path ${expectationURL.pathname}
-Expectation should be of the form path/to/cts.https.html?worker=0&q=suite:test_path:test_name:foo=1;bar=2;...
+Expectation should be of the form path/to/cts.https.html?debug=0&q=suite:test_path:test_name:foo=1;bar=2;...
         `
       );
 
       const params = expectationURL.searchParams;
-      if (optionEnabled('worker', params) !== optionEnabled('worker', wptURL.searchParams)) {
+      if (optionString('worker', params) !== optionString('worker', wptURL.searchParams)) {
         continue;
       }
 
diff --git a/src/common/runtime/helper/options.ts b/src/common/runtime/helper/options.ts
index 80bd4bccc2f4..60c2d1069181 100644
--- a/src/common/runtime/helper/options.ts
+++ b/src/common/runtime/helper/options.ts
@@ -25,7 +25,7 @@ export function optionString(
  * The possible options for the tests.
  */
 export interface CTSOptions {
-  worker: boolean;
+  worker?: 'dedicated' | 'shared' | '';
   debug: boolean;
   compatibility: boolean;
   forceFallbackAdapter: boolean;
@@ -34,7 +34,7 @@ export interface CTSOptions {
 }
 
 export const kDefaultCTSOptions: CTSOptions = {
-  worker: false,
+  worker: '',
   debug: true,
   compatibility: false,
   forceFallbackAdapter: false,
@@ -61,7 +61,15 @@ export type OptionsInfos<Type> = Record<keyof Type, OptionInfo>;
  * Options to the CTS.
  */
 export const kCTSOptionsInfo: OptionsInfos<CTSOptions> = {
-  worker: { description: 'run in a worker' },
+  worker: {
+    description: 'run in a worker',
+    parser: optionString,
+    selectValueDescriptions: [
+      { value: '', description: 'no worker' },
+      { value: 'dedicated', description: 'dedicated worker' },
+      { value: 'shared', description: 'shared worker' },
+    ],
+  },
   debug: { description: 'show more info' },
   compatibility: { description: 'run in compatibility mode' },
   forceFallbackAdapter: { description: 'pass forceFallbackAdapter: true to requestAdapter' },
diff --git a/src/common/runtime/helper/test_worker-worker.ts b/src/common/runtime/helper/test_worker-worker.ts
index e8d187ea7e3d..caf6e7a1bc17 100644
--- a/src/common/runtime/helper/test_worker-worker.ts
+++ b/src/common/runtime/helper/test_worker-worker.ts
@@ -9,7 +9,7 @@ import { assert } from '../../util/util.js';
 
 import { CTSOptions } from './options.js';
 
-// Should be DedicatedWorkerGlobalScope, but importing lib "webworker" conflicts with lib "dom".
+// Should be WorkerGlobalScope, but importing lib "webworker" conflicts with lib "dom".
 /* eslint-disable-next-line @typescript-eslint/no-explicit-any */
 declare const self: any;
 
@@ -17,7 +17,7 @@ const loader = new DefaultTestFileLoader();
 
 setBaseResourcePath('../../../resources');
 
-self.onmessage = async (ev: MessageEvent) => {
+async function reportTestResults(this: MessagePort | Worker, ev: MessageEvent) {
   const query: string = ev.data.query;
   const expectations: TestQueryWithExpectation[] = ev.data.expectations;
   const ctsOptions: CTSOptions = ev.data.ctsOptions;
@@ -44,5 +44,17 @@ self.onmessage = async (ev: MessageEvent) => {
   const [rec, result] = log.record(testcase.query.toString());
   await testcase.run(rec, expectations);
 
-  self.postMessage({ query, result });
+  this.postMessage({ query, result });
+}
+
+self.onmessage = (ev: MessageEvent) => {
+  void reportTestResults.call(self, ev);
+};
+
+self.onconnect = (event: MessageEvent) => {
+  const port = event.ports[0];
+
+  port.onmessage = (ev: MessageEvent) => {
+    void reportTestResults.call(port, ev);
+  };
 };
diff --git a/src/common/runtime/helper/test_worker.ts b/src/common/runtime/helper/test_worker.ts
index 9bbcab094671..a9ca4e2338c7 100644
--- a/src/common/runtime/helper/test_worker.ts
+++ b/src/common/runtime/helper/test_worker.ts
@@ -5,13 +5,13 @@ import { TestQueryWithExpectation } from '../../internal/query/query.js';
 
 import { CTSOptions, kDefaultCTSOptions } from './options.js';
 
-export class TestWorker {
+export class TestDedicatedWorker {
   private readonly ctsOptions: CTSOptions;
   private readonly worker: Worker;
   private readonly resolvers = new Map<string, (result: LiveTestCaseResult) => void>();
 
   constructor(ctsOptions?: CTSOptions) {
-    this.ctsOptions = { ...(ctsOptions || kDefaultCTSOptions), ...{ worker: true } };
+    this.ctsOptions = { ...(ctsOptions || kDefaultCTSOptions), ...{ worker: 'dedicated' } };
     const selfPath = import.meta.url;
     const selfPathDir = selfPath.substring(0, selfPath.lastIndexOf('/'));
     const workerPath = selfPathDir + '/test_worker-worker.js';
@@ -47,3 +47,48 @@ export class TestWorker {
     rec.injectResult(workerResult);
   }
 }
+
+export class TestSharedWorker {
+  private readonly ctsOptions: CTSOptions;
+  private readonly port: MessagePort;
+  private readonly resolvers = new Map<string, (result: LiveTestCaseResult) => void>();
+
+  constructor(ctsOptions?: CTSOptions) {
+    this.ctsOptions = { ...(ctsOptions || kDefaultCTSOptions), ...{ worker: 'shared' } };
+    const selfPath = import.meta.url;
+    const selfPathDir = selfPath.substring(0, selfPath.lastIndexOf('/'));
+    const workerPath = selfPathDir + '/test_worker-worker.js';
+    const worker = new SharedWorker(workerPath, { type: 'module' });
+    this.port = worker.port;
+    this.port.start();
+    this.port.onmessage = ev => {
+      const query: string = ev.data.query;
+      const result: TransferredTestCaseResult = ev.data.result;
+      if (result.logs) {
+        for (const l of result.logs) {
+          Object.setPrototypeOf(l, LogMessageWithStack.prototype);
+        }
+      }
+      this.resolvers.get(query)!(result as LiveTestCaseResult);
+
+      // MAINTENANCE_TODO(kainino0x): update the Logger with this result (or don't have a logger and
+      // update the entire results JSON somehow at some point).
+    };
+  }
+
+  async run(
+    rec: TestCaseRecorder,
+    query: string,
+    expectations: TestQueryWithExpectation[] = []
+  ): Promise<void> {
+    this.port.postMessage({
+      query,
+      expectations,
+      ctsOptions: this.ctsOptions,
+    });
+    const workerResult = await new Promise<LiveTestCaseResult>(resolve => {
+      this.resolvers.set(query, resolve);
+    });
+    rec.injectResult(workerResult);
+  }
+}
diff --git a/src/common/runtime/standalone.ts b/src/common/runtime/standalone.ts
index aaee30df6c5e..fa938c55a276 100644
--- a/src/common/runtime/standalone.ts
+++ b/src/common/runtime/standalone.ts
@@ -21,7 +21,7 @@ import {
   OptionsInfos,
   camelCaseToSnakeCase,
 } from './helper/options.js';
-import { TestWorker } from './helper/test_worker.js';
+import { TestDedicatedWorker, TestSharedWorker } from './helper/test_worker.js';
 
 const rootQuerySpec = 'webgpu:*';
 let promptBeforeReload = false;
@@ -63,7 +63,9 @@ const logger = new Logger();
 
 setBaseResourcePath('../out/resources');
 
-const worker = options.worker ? new TestWorker(options) : undefined;
+const dedicatedWorker =
+  options.worker === 'dedicated' ? new TestDedicatedWorker(options) : undefined;
+const sharedWorker = options.worker === 'shared' ? new TestSharedWorker(options) : undefined;
 
 const autoCloseOnPass = document.getElementById('autoCloseOnPass') as HTMLInputElement;
 const resultsVis = document.getElementById('resultsVis')!;
@@ -176,8 +178,10 @@ function makeCaseHTML(t: TestTreeLeaf): VisualizedSubtree {
 
     const [rec, res] = logger.record(name);
     caseResult = res;
-    if (worker) {
-      await worker.run(rec, name);
+    if (dedicatedWorker) {
+      await dedicatedWorker.run(rec, name);
+    } else if (sharedWorker) {
+      await sharedWorker.run(rec, name);
     } else {
       await t.run(rec);
     }
diff --git a/src/common/runtime/wpt.ts b/src/common/runtime/wpt.ts
index d4a40081540e..aacd34b13e51 100644
--- a/src/common/runtime/wpt.ts
+++ b/src/common/runtime/wpt.ts
@@ -8,8 +8,8 @@ import { parseQuery } from '../internal/query/parseQuery.js';
 import { parseExpectationsForTestQuery, relativeQueryString } from '../internal/query/query.js';
 import { assert } from '../util/util.js';
 
-import { optionEnabled } from './helper/options.js';
-import { TestWorker } from './helper/test_worker.js';
+import { optionEnabled, optionString } from './helper/options.js';
+import { TestDedicatedWorker, TestSharedWorker } from './helper/test_worker.js';
 
 // testharness.js API (https://web-platform-tests.org/writing-tests/testharness-api.html)
 declare interface WptTestObject {
@@ -31,8 +31,9 @@ setup({
 });
 
 void (async () => {
-  const workerEnabled = optionEnabled('worker');
-  const worker = workerEnabled ? new TestWorker() : undefined;
+  const workerString = optionString('worker');
+  const dedicatedWorker = workerString === 'dedicated' ? new TestDedicatedWorker() : undefined;
+  const sharedWorker = workerString === 'shared' ? new TestSharedWorker() : undefined;
 
   globalTestConfig.unrollConstEvalLoops = optionEnabled('unroll_const_eval_loops');
 
@@ -63,8 +64,10 @@ void (async () => {
 
     const wpt_fn = async () => {
       const [rec, res] = log.record(name);
-      if (worker) {
-        await worker.run(rec, name, expectations);
+      if (dedicatedWorker) {
+        await dedicatedWorker.run(rec, name, expectations);
+      } else if (sharedWorker) {
+        await sharedWorker.run(rec, name, expectations);
       } else {
         await testcase.run(rec, expectations);
       }
diff --git a/src/common/tools/gen_wpt_cts_html.ts b/src/common/tools/gen_wpt_cts_html.ts
index 90c7cd4ef4c5..8293d6612671 100644
--- a/src/common/tools/gen_wpt_cts_html.ts
+++ b/src/common/tools/gen_wpt_cts_html.ts
@@ -35,15 +35,15 @@ where arguments.txt is a file containing a list of arguments prefixes to both ge
 in the expectations. The entire variant list generation runs *once per prefix*, so this
 multiplies the size of the variant list.
 
-  ?worker=0&q=
-  ?worker=1&q=
+  ?debug=0&q=
+  ?debug=1&q=
 
 and myexpectations.txt is a file containing a list of WPT paths to suppress, e.g.:
 
-  path/to/cts.https.html?worker=0&q=webgpu:a/foo:bar={"x":1}
-  path/to/cts.https.html?worker=1&q=webgpu:a/foo:bar={"x":1}
+  path/to/cts.https.html?debug=0&q=webgpu:a/foo:bar={"x":1}
+  path/to/cts.https.html?debug=1&q=webgpu:a/foo:bar={"x":1}
 
-  path/to/cts.https.html?worker=1&q=webgpu:a/foo:bar={"x":3}
+  path/to/cts.https.html?debug=1&q=webgpu:a/foo:bar={"x":3}
 `);
   process.exit(rc);
 }
@@ -224,7 +224,7 @@ ${queryString}`
       }
 
       lines.push({
-        urlQueryString: prefix + query.toString(), // "?worker=0&q=..."
+        urlQueryString: prefix + query.toString(), // "?debug=0&q=..."
         comment: useChunking ? `estimated: ${subtreeCounts?.totalTimeMS.toFixed(3)} ms` : undefined,
       });
 
diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json
index 791bd3a0354b..73df50c4d73f 100644
--- a/src/resources/cache/hashes.json
+++ b/src/resources/cache/hashes.json
@@ -1,109 +1,109 @@
 {
-  "webgpu/shader/execution/binary/af_addition.bin": "6a22792d",
-  "webgpu/shader/execution/binary/af_logical.bin": "d124cf64",
-  "webgpu/shader/execution/binary/af_division.bin": "7d78cca7",
-  "webgpu/shader/execution/binary/af_matrix_addition.bin": "fd87a833",
-  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "814fc609",
-  "webgpu/shader/execution/binary/af_multiplication.bin": "5a2c3c94",
-  "webgpu/shader/execution/binary/af_remainder.bin": "f1a63570",
-  "webgpu/shader/execution/binary/af_subtraction.bin": "3ecadf73",
-  "webgpu/shader/execution/binary/f16_addition.bin": "6c64f032",
-  "webgpu/shader/execution/binary/f16_logical.bin": "3120fe75",
-  "webgpu/shader/execution/binary/f16_division.bin": "a432967b",
-  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "42c4a69e",
-  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "9a04c20",
-  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "8c4dcf3c",
-  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "755f369d",
-  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "eae889e3",
-  "webgpu/shader/execution/binary/f16_multiplication.bin": "53d2e0ea",
-  "webgpu/shader/execution/binary/f16_remainder.bin": "c0f8829e",
-  "webgpu/shader/execution/binary/f16_subtraction.bin": "d63e8b7b",
-  "webgpu/shader/execution/binary/f32_addition.bin": "7bd51700",
-  "webgpu/shader/execution/binary/f32_logical.bin": "72dfd230",
-  "webgpu/shader/execution/binary/f32_division.bin": "c93dbaa5",
-  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "bb08a2c6",
-  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "e61f22cf",
-  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "a8c61bd4",
-  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "82d727d",
-  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "1c21512",
-  "webgpu/shader/execution/binary/f32_multiplication.bin": "dadd324d",
-  "webgpu/shader/execution/binary/f32_remainder.bin": "6a3bac80",
-  "webgpu/shader/execution/binary/f32_subtraction.bin": "813f8e0d",
-  "webgpu/shader/execution/binary/i32_arithmetic.bin": "12b039ce",
-  "webgpu/shader/execution/binary/i32_comparison.bin": "dec54fb8",
-  "webgpu/shader/execution/binary/u32_arithmetic.bin": "548fc8ea",
-  "webgpu/shader/execution/binary/u32_comparison.bin": "62be0c6d",
-  "webgpu/shader/execution/abs.bin": "cf195671",
-  "webgpu/shader/execution/acos.bin": "7b18783a",
-  "webgpu/shader/execution/acosh.bin": "1f24550",
-  "webgpu/shader/execution/asin.bin": "9afbbed6",
-  "webgpu/shader/execution/asinh.bin": "e1c343ec",
-  "webgpu/shader/execution/atan.bin": "8457e2a8",
-  "webgpu/shader/execution/atan2.bin": "f20f1292",
-  "webgpu/shader/execution/atanh.bin": "dc1a948b",
-  "webgpu/shader/execution/bitcast.bin": "67023663",
-  "webgpu/shader/execution/ceil.bin": "f98273c3",
-  "webgpu/shader/execution/clamp.bin": "fb634e90",
-  "webgpu/shader/execution/cos.bin": "33cce38a",
-  "webgpu/shader/execution/cosh.bin": "1c55ac51",
-  "webgpu/shader/execution/cross.bin": "d5a28621",
-  "webgpu/shader/execution/degrees.bin": "4fcc3175",
-  "webgpu/shader/execution/determinant.bin": "4e335eb0",
-  "webgpu/shader/execution/distance.bin": "d39f063c",
-  "webgpu/shader/execution/dot.bin": "4feda0d8",
-  "webgpu/shader/execution/exp.bin": "4c948963",
-  "webgpu/shader/execution/exp2.bin": "454497e2",
-  "webgpu/shader/execution/faceForward.bin": "97413daf",
-  "webgpu/shader/execution/floor.bin": "7da41e10",
-  "webgpu/shader/execution/fma.bin": "65588609",
-  "webgpu/shader/execution/fract.bin": "47272979",
-  "webgpu/shader/execution/frexp.bin": "50c7d61a",
-  "webgpu/shader/execution/inverseSqrt.bin": "50261c42",
-  "webgpu/shader/execution/ldexp.bin": "1fb40fae",
-  "webgpu/shader/execution/length.bin": "738be1d5",
-  "webgpu/shader/execution/log.bin": "9c5e72c3",
-  "webgpu/shader/execution/log2.bin": "b2a40016",
-  "webgpu/shader/execution/max.bin": "85e864a8",
-  "webgpu/shader/execution/min.bin": "f9a02244",
-  "webgpu/shader/execution/mix.bin": "2b126c4f",
-  "webgpu/shader/execution/modf.bin": "6321456f",
-  "webgpu/shader/execution/normalize.bin": "5e1fab03",
-  "webgpu/shader/execution/pack2x16float.bin": "a0ee7786",
-  "webgpu/shader/execution/pow.bin": "94dbc1e8",
-  "webgpu/shader/execution/quantizeToF16.bin": "f9072ec0",
-  "webgpu/shader/execution/radians.bin": "bce17911",
-  "webgpu/shader/execution/reflect.bin": "f9f4128a",
-  "webgpu/shader/execution/refract.bin": "a1829a6e",
-  "webgpu/shader/execution/round.bin": "4e431ad7",
-  "webgpu/shader/execution/saturate.bin": "dea37ddb",
-  "webgpu/shader/execution/sign.bin": "830ec10",
-  "webgpu/shader/execution/sin.bin": "922531b5",
-  "webgpu/shader/execution/sinh.bin": "3193d989",
-  "webgpu/shader/execution/smoothstep.bin": "51503cb",
-  "webgpu/shader/execution/sqrt.bin": "f933e0d2",
-  "webgpu/shader/execution/step.bin": "ce02367d",
-  "webgpu/shader/execution/tan.bin": "b9a3a310",
-  "webgpu/shader/execution/tanh.bin": "98685dd3",
-  "webgpu/shader/execution/transpose.bin": "bef30077",
-  "webgpu/shader/execution/trunc.bin": "b8b5c091",
-  "webgpu/shader/execution/unpack2x16float.bin": "9343b29f",
-  "webgpu/shader/execution/unpack2x16snorm.bin": "e56ec75e",
-  "webgpu/shader/execution/unpack2x16unorm.bin": "26504819",
-  "webgpu/shader/execution/unpack4x8snorm.bin": "e550de72",
-  "webgpu/shader/execution/unpack4x8unorm.bin": "acdf694f",
-  "webgpu/shader/execution/unary/af_arithmetic.bin": "d371eefd",
-  "webgpu/shader/execution/unary/af_assignment.bin": "4cae15c1",
-  "webgpu/shader/execution/unary/bool_conversion.bin": "be4a17e6",
-  "webgpu/shader/execution/unary/f16_arithmetic.bin": "cf8bd6e9",
-  "webgpu/shader/execution/unary/f16_conversion.bin": "44ea8241",
-  "webgpu/shader/execution/unary/f32_arithmetic.bin": "e046d177",
-  "webgpu/shader/execution/unary/f32_conversion.bin": "9ef137d6",
-  "webgpu/shader/execution/unary/i32_arithmetic.bin": "44c3fd63",
-  "webgpu/shader/execution/unary/i32_complement.bin": "82e58655",
-  "webgpu/shader/execution/unary/i32_conversion.bin": "4b164846",
-  "webgpu/shader/execution/unary/u32_complement.bin": "51de4a0f",
-  "webgpu/shader/execution/unary/u32_conversion.bin": "659c2f08",
-  "webgpu/shader/execution/unary/ai_assignment.bin": "bf7d58b9",
-  "webgpu/shader/execution/binary/ai_arithmetic.bin": "7881600d",
-  "webgpu/shader/execution/unary/ai_arithmetic.bin": "eda6e030"
+  "webgpu/shader/execution/binary/af_addition.bin": "f9c25825",
+  "webgpu/shader/execution/binary/af_logical.bin": "c3c5c046",
+  "webgpu/shader/execution/binary/af_division.bin": "8eec7360",
+  "webgpu/shader/execution/binary/af_matrix_addition.bin": "92f5acd",
+  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "cf1920bc",
+  "webgpu/shader/execution/binary/af_multiplication.bin": "c3ca2d36",
+  "webgpu/shader/execution/binary/af_remainder.bin": "ef236c3e",
+  "webgpu/shader/execution/binary/af_subtraction.bin": "823c3d19",
+  "webgpu/shader/execution/binary/f16_addition.bin": "72625726",
+  "webgpu/shader/execution/binary/f16_logical.bin": "d20f2684",
+  "webgpu/shader/execution/binary/f16_division.bin": "865fb316",
+  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "2908c1e6",
+  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "dbc6dd4a",
+  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "fff001f3",
+  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "f6557db8",
+  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "7d0beef1",
+  "webgpu/shader/execution/binary/f16_multiplication.bin": "fe1eeb2f",
+  "webgpu/shader/execution/binary/f16_remainder.bin": "44c68db0",
+  "webgpu/shader/execution/binary/f16_subtraction.bin": "4d349e03",
+  "webgpu/shader/execution/binary/f32_addition.bin": "3d909f18",
+  "webgpu/shader/execution/binary/f32_logical.bin": "300cb917",
+  "webgpu/shader/execution/binary/f32_division.bin": "26126d0b",
+  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "8e3b7347",
+  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "eba8334f",
+  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "5f633cc4",
+  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "ee699077",
+  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "6e8e9044",
+  "webgpu/shader/execution/binary/f32_multiplication.bin": "9e8aa842",
+  "webgpu/shader/execution/binary/f32_remainder.bin": "44997694",
+  "webgpu/shader/execution/binary/f32_subtraction.bin": "7da0351c",
+  "webgpu/shader/execution/binary/i32_arithmetic.bin": "78991d8",
+  "webgpu/shader/execution/binary/i32_comparison.bin": "eeb5a624",
+  "webgpu/shader/execution/binary/u32_arithmetic.bin": "e317f38",
+  "webgpu/shader/execution/binary/u32_comparison.bin": "3bd52719",
+  "webgpu/shader/execution/abs.bin": "58d33fc0",
+  "webgpu/shader/execution/acos.bin": "8ab770e1",
+  "webgpu/shader/execution/acosh.bin": "1b797a55",
+  "webgpu/shader/execution/asin.bin": "49c3e54",
+  "webgpu/shader/execution/asinh.bin": "ec965820",
+  "webgpu/shader/execution/atan.bin": "13292f0c",
+  "webgpu/shader/execution/atan2.bin": "943937cf",
+  "webgpu/shader/execution/atanh.bin": "113723aa",
+  "webgpu/shader/execution/bitcast.bin": "49055551",
+  "webgpu/shader/execution/ceil.bin": "fc4d440",
+  "webgpu/shader/execution/clamp.bin": "f0590b61",
+  "webgpu/shader/execution/cos.bin": "dc6a6d3c",
+  "webgpu/shader/execution/cosh.bin": "682a8b93",
+  "webgpu/shader/execution/cross.bin": "c2cf537f",
+  "webgpu/shader/execution/degrees.bin": "b9044d85",
+  "webgpu/shader/execution/determinant.bin": "c872e66b",
+  "webgpu/shader/execution/distance.bin": "8df8d8a7",
+  "webgpu/shader/execution/dot.bin": "84aacdd9",
+  "webgpu/shader/execution/exp.bin": "1e815ae0",
+  "webgpu/shader/execution/exp2.bin": "ee9fd571",
+  "webgpu/shader/execution/faceForward.bin": "e2b2e81c",
+  "webgpu/shader/execution/floor.bin": "9d437a66",
+  "webgpu/shader/execution/fma.bin": "352b140",
+  "webgpu/shader/execution/fract.bin": "95a1650d",
+  "webgpu/shader/execution/frexp.bin": "c9ed5769",
+  "webgpu/shader/execution/inverseSqrt.bin": "8a275a8d",
+  "webgpu/shader/execution/ldexp.bin": "cb51dbe1",
+  "webgpu/shader/execution/length.bin": "53468ea4",
+  "webgpu/shader/execution/log.bin": "33bef86d",
+  "webgpu/shader/execution/log2.bin": "42fc6f07",
+  "webgpu/shader/execution/max.bin": "33b63886",
+  "webgpu/shader/execution/min.bin": "f4e85c7",
+  "webgpu/shader/execution/mix.bin": "b7346332",
+  "webgpu/shader/execution/modf.bin": "b291939c",
+  "webgpu/shader/execution/normalize.bin": "8d42f4e4",
+  "webgpu/shader/execution/pack2x16float.bin": "32836353",
+  "webgpu/shader/execution/pow.bin": "5fd2851d",
+  "webgpu/shader/execution/quantizeToF16.bin": "7a1878ee",
+  "webgpu/shader/execution/radians.bin": "dbb58186",
+  "webgpu/shader/execution/reflect.bin": "9f1132c5",
+  "webgpu/shader/execution/refract.bin": "508a9b82",
+  "webgpu/shader/execution/round.bin": "47ccb377",
+  "webgpu/shader/execution/saturate.bin": "771426ff",
+  "webgpu/shader/execution/sign.bin": "6b04f125",
+  "webgpu/shader/execution/sin.bin": "8e2495ec",
+  "webgpu/shader/execution/sinh.bin": "77ab8cf1",
+  "webgpu/shader/execution/smoothstep.bin": "bdd40525",
+  "webgpu/shader/execution/sqrt.bin": "d6b16a43",
+  "webgpu/shader/execution/step.bin": "34fd45ba",
+  "webgpu/shader/execution/tan.bin": "7df2e6b4",
+  "webgpu/shader/execution/tanh.bin": "e303ee3c",
+  "webgpu/shader/execution/transpose.bin": "985b9dd0",
+  "webgpu/shader/execution/trunc.bin": "b244b3d9",
+  "webgpu/shader/execution/unpack2x16float.bin": "c1ebb8d5",
+  "webgpu/shader/execution/unpack2x16snorm.bin": "f7fd89e",
+  "webgpu/shader/execution/unpack2x16unorm.bin": "35e0330e",
+  "webgpu/shader/execution/unpack4x8snorm.bin": "3d4ca4cd",
+  "webgpu/shader/execution/unpack4x8unorm.bin": "985e1f9a",
+  "webgpu/shader/execution/unary/af_arithmetic.bin": "fab31dec",
+  "webgpu/shader/execution/unary/af_assignment.bin": "af1899",
+  "webgpu/shader/execution/unary/bool_conversion.bin": "c4e6e40c",
+  "webgpu/shader/execution/unary/f16_arithmetic.bin": "db964478",
+  "webgpu/shader/execution/unary/f16_conversion.bin": "df65cc9a",
+  "webgpu/shader/execution/unary/f32_arithmetic.bin": "bd022e24",
+  "webgpu/shader/execution/unary/f32_conversion.bin": "ff00ace3",
+  "webgpu/shader/execution/unary/i32_arithmetic.bin": "239e3b46",
+  "webgpu/shader/execution/unary/i32_complement.bin": "81fe7e5b",
+  "webgpu/shader/execution/unary/i32_conversion.bin": "24deb50a",
+  "webgpu/shader/execution/unary/u32_complement.bin": "37ba2547",
+  "webgpu/shader/execution/unary/u32_conversion.bin": "f26f8c3a",
+  "webgpu/shader/execution/unary/ai_assignment.bin": "46ee1c6f",
+  "webgpu/shader/execution/binary/ai_arithmetic.bin": "aa09a4bc",
+  "webgpu/shader/execution/unary/ai_arithmetic.bin": "6e2d927f"
 }
\ No newline at end of file
diff --git a/src/resources/cache/webgpu/shader/execution/bitcast.bin b/src/resources/cache/webgpu/shader/execution/bitcast.bin
index 3c24c88e1ffbc42f87639912bacab411154e6d63..ab435f3902a6e2124845cfba3625a2184be02b02 100644
GIT binary patch
delta 152
zcmWN=yAgsw6hP5g1Ox<KK=~*Nl)x*{gkU$?ggzL?#3Jm7$a`jv{TXGsy=7S^Jv%W@
u6g?j3GhoOgBgRa4V#<tXUZ|O~V96_QtXQ*Q%R3+J_~eT{hp&@y|NH?yCN>ZN

delta 151
zcmWN=#|eNy002SEIp?e`T!C`nL(36-ScKLj?!X++t{gm9p}yZH#V&<Hl^S&#G-=VM
jLzf<X1`HW7X2O&ia~3RFv1Y@T9eWNOIdQ(FSnTTu*p)Ml

diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index 9826a3e806ab..c6647081079d 100644
--- a/src/webgpu/listing_meta.json
+++ b/src/webgpu/listing_meta.json
@@ -609,7 +609,7 @@
   "webgpu:api,validation,image_copy,texture_related:texture,device_mismatch:*": { "subcaseMS": 5.417 },
   "webgpu:api,validation,image_copy,texture_related:usage:*": { "subcaseMS": 1.224 },
   "webgpu:api,validation,image_copy,texture_related:valid:*": { "subcaseMS": 3.678 },
-  "webgpu:api,validation,layout_shader_compat:pipeline_layout_shader_exact_match:*": { "subcaseMS": 2.00 },
+  "webgpu:api,validation,layout_shader_compat:pipeline_layout_shader_exact_match:*": { "subcaseMS": 2.000 },
   "webgpu:api,validation,query_set,create:count:*": { "subcaseMS": 0.967 },
   "webgpu:api,validation,query_set,destroy:invalid_queryset:*": { "subcaseMS": 0.801 },
   "webgpu:api,validation,query_set,destroy:twice:*": { "subcaseMS": 0.700 },
@@ -674,13 +674,13 @@
   "webgpu:api,validation,render_pass,render_pass_descriptor:attachments,one_color_attachment:*": { "subcaseMS": 33.401 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:attachments,one_depth_stencil_attachment:*": { "subcaseMS": 15.301 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:attachments,same_size:*": { "subcaseMS": 33.400 },
+  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,bound_check:*": { "subcaseMS": 9.400 },
+  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,definedness:*": { "subcaseMS": 5.601 },
+  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,overlaps,diff_miplevel:*": { "subcaseMS": 3.901 },
+  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,overlaps,same_miplevel:*": { "subcaseMS": 6.400 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,empty:*": { "subcaseMS": 0.400 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,limits,maxColorAttachmentBytesPerSample,aligned:*": { "subcaseMS": 1.825 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,limits,maxColorAttachmentBytesPerSample,unaligned:*": { "subcaseMS": 17.151 },
-  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,definedness:*": { "subcaseMS": 5.601 },
-  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,bound_check:*": { "subcaseMS": 9.400 },
-  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,overlaps,same_miplevel:*": { "subcaseMS": 6.400 },
-  "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,depthSlice,overlaps,diff_miplevel:*": { "subcaseMS": 3.901 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,limits,maxColorAttachments:*": { "subcaseMS": 0.950 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,non_multisampled:*": { "subcaseMS": 32.601 },
   "webgpu:api,validation,render_pass,render_pass_descriptor:color_attachments,sample_count:*": { "subcaseMS": 33.600 },
@@ -897,21 +897,6 @@
   "webgpu:idl,constructable:gpu_errors:*": { "subcaseMS": 0.101 },
   "webgpu:idl,constructable:pipeline_errors:*": { "subcaseMS": 0.101 },
   "webgpu:idl,constructable:uncaptured_error_event:*": { "subcaseMS": 0.101 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition_scalar_vector:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition_vector_scalar:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:division:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:division_scalar_vector:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:division_vector_scalar:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication_scalar_vector:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder_scalar_vector:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder_vector_scalar:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication_vector_scalar:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction_scalar_vector:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction_vector_scalar:*": { "subcaseMS": 0 },
   "webgpu:shader,execution,expression,binary,af_addition:scalar:*": { "subcaseMS": 815.300 },
   "webgpu:shader,execution,expression,binary,af_addition:scalar_vector:*": { "subcaseMS": 1803.434 },
   "webgpu:shader,execution,expression,binary,af_addition:vector:*": { "subcaseMS": 719.600 },
@@ -940,6 +925,21 @@
   "webgpu:shader,execution,expression,binary,af_subtraction:scalar_vector:*": { "subcaseMS": 2336.534 },
   "webgpu:shader,execution,expression,binary,af_subtraction:vector:*": { "subcaseMS": 764.201 },
   "webgpu:shader,execution,expression,binary,af_subtraction:vector_scalar:*": { "subcaseMS": 2437.701 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition_scalar_vector:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:addition_vector_scalar:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:division:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:division_scalar_vector:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:division_vector_scalar:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication_scalar_vector:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:multiplication_vector_scalar:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder_scalar_vector:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:remainder_vector_scalar:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction_scalar_vector:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,binary,ai_arithmetic:subtraction_vector_scalar:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,binary,bitwise:bitwise_and:*": { "subcaseMS": 20.982 },
   "webgpu:shader,execution,expression,binary,bitwise:bitwise_and_compound:*": { "subcaseMS": 22.513 },
   "webgpu:shader,execution,expression,binary,bitwise:bitwise_exclusive_or:*": { "subcaseMS": 21.294 },
@@ -1178,6 +1178,14 @@
   "webgpu:shader,execution,expression,call,builtin,atomics,atomicSub:sub_workgroup:*": { "subcaseMS": 7.238 },
   "webgpu:shader,execution,expression,call,builtin,atomics,atomicXor:xor_storage:*": { "subcaseMS": 6.807 },
   "webgpu:shader,execution,expression,call,builtin,atomics,atomicXor:xor_workgroup:*": { "subcaseMS": 7.821 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_f32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_i32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_u32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_vec2f16:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_f32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_i32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_u32:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_vec2f16:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:f16_to_f16:*": { "subcaseMS": 21.112 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:f32_to_f32:*": { "subcaseMS": 8.625 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:f32_to_i32:*": { "subcaseMS": 8.175 },
@@ -1191,6 +1199,8 @@
   "webgpu:shader,execution,expression,call,builtin,bitcast:u32_to_i32:*": { "subcaseMS": 6.982 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:u32_to_u32:*": { "subcaseMS": 6.907 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:u32_to_vec2h:*": { "subcaseMS": 22.210 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:vec2af_to_vec4f16:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,builtin,bitcast:vec2ai_to_vec4f16:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:vec2f_to_vec4h:*": { "subcaseMS": 24.015 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:vec2h_to_f32:*": { "subcaseMS": 21.412 },
   "webgpu:shader,execution,expression,call,builtin,bitcast:vec2h_to_i32:*": { "subcaseMS": 38.312 },
@@ -1239,6 +1249,8 @@
   "webgpu:shader,execution,expression,call,builtin,distance:f32_vec2:*": { "subcaseMS": 9.826 },
   "webgpu:shader,execution,expression,call,builtin,distance:f32_vec3:*": { "subcaseMS": 10.901 },
   "webgpu:shader,execution,expression,call,builtin,distance:f32_vec4:*": { "subcaseMS": 12.700 },
+  "webgpu:shader,execution,expression,call,builtin,dot4I8Packed:basic:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,execution,expression,call,builtin,dot4U8Packed:basic:*": { "subcaseMS": 1.000 },
   "webgpu:shader,execution,expression,call,builtin,dot:abstract_float:*": { "subcaseMS": 8.902 },
   "webgpu:shader,execution,expression,call,builtin,dot:abstract_int:*": { "subcaseMS": 2.902 },
   "webgpu:shader,execution,expression,call,builtin,dot:f16_vec2:*": { "subcaseMS": 981.225 },
@@ -1249,8 +1261,6 @@
   "webgpu:shader,execution,expression,call,builtin,dot:f32_vec4:*": { "subcaseMS": 11.876 },
   "webgpu:shader,execution,expression,call,builtin,dot:i32:*": { "subcaseMS": 3.103 },
   "webgpu:shader,execution,expression,call,builtin,dot:u32:*": { "subcaseMS": 3.101 },
-  "webgpu:shader,execution,expression,call,builtin,dot4I8Packed:basic:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,execution,expression,call,builtin,dot4U8Packed:basic:*": { "subcaseMS": 1.000 },
   "webgpu:shader,execution,expression,call,builtin,dpdx:f32:*": { "subcaseMS": 22.804 },
   "webgpu:shader,execution,expression,call,builtin,dpdxCoarse:f32:*": { "subcaseMS": 22.404 },
   "webgpu:shader,execution,expression,call,builtin,dpdxFine:f32:*": { "subcaseMS": 17.708 },
@@ -1538,18 +1548,21 @@
   "webgpu:shader,execution,expression,call,builtin,unpack4xU8:basic:*": { "subcaseMS": 78.501 },
   "webgpu:shader,execution,expression,call,builtin,workgroupBarrier:barrier:*": { "subcaseMS": 0.701 },
   "webgpu:shader,execution,expression,call,builtin,workgroupBarrier:stage:*": { "subcaseMS": 1.801 },
+  "webgpu:shader,execution,expression,call,user,ptr_params:mixed_ptr_parameters:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,user,ptr_params:read_full_object:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,execution,expression,call,user,ptr_params:read_ptr_to_member:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,user,ptr_params:read_ptr_to_element:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,user,ptr_params:read_ptr_to_member:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,user,ptr_params:write_full_object:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,execution,expression,call,user,ptr_params:write_ptr_to_member:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,call,user,ptr_params:write_ptr_to_element:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,execution,expression,call,user,ptr_params:mixed_ptr_parameters:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,execution,expression,unary,ai_arithmetic:negation:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,call,user,ptr_params:write_ptr_to_member:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_index:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_member:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,unary,af_arithmetic:negation:*": { "subcaseMS": 2165.950 },
   "webgpu:shader,execution,expression,unary,af_assignment:abstract:*": { "subcaseMS": 788.400 },
   "webgpu:shader,execution,expression,unary,af_assignment:f16:*": { "subcaseMS": 1.000 },
   "webgpu:shader,execution,expression,unary,af_assignment:f32:*": { "subcaseMS": 42.000 },
+  "webgpu:shader,execution,expression,unary,ai_arithmetic:negation:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,unary,ai_assignment:abstract:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,unary,ai_assignment:i32:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,expression,unary,ai_assignment:u32:*": { "subcaseMS": 0.000 },
@@ -1559,9 +1572,6 @@
   "webgpu:shader,execution,expression,unary,bool_conversion:i32:*": { "subcaseMS": 8.219 },
   "webgpu:shader,execution,expression,unary,bool_conversion:u32:*": { "subcaseMS": 7.401 },
   "webgpu:shader,execution,expression,unary,bool_logical:negation:*": { "subcaseMS": 6.413 },
-  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref:*": { "subcaseMS": 0.0 },
-  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_index:*": { "subcaseMS": 0.0 },
-  "webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_member:*": { "subcaseMS": 0.0 },
   "webgpu:shader,execution,expression,unary,f16_arithmetic:negation:*": { "subcaseMS": 117.604 },
   "webgpu:shader,execution,expression,unary,f16_conversion:bool:*": { "subcaseMS": 34.694 },
   "webgpu:shader,execution,expression,unary,f16_conversion:f16:*": { "subcaseMS": 36.013 },
@@ -1670,8 +1680,8 @@
   "webgpu:shader,execution,flow_control,while:while_continue:*": { "subcaseMS": 5.650 },
   "webgpu:shader,execution,flow_control,while:while_nested_break:*": { "subcaseMS": 12.701 },
   "webgpu:shader,execution,flow_control,while:while_nested_continue:*": { "subcaseMS": 5.450 },
-  "webgpu:shader,execution,memory_layout:read_layout:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,memory_layout:write_layout:*": { "subcaseMS": 0 },
+  "webgpu:shader,execution,memory_layout:read_layout:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,execution,memory_layout:write_layout:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,memory_model,adjacent:f16:*": { "subcaseMS": 23.625 },
   "webgpu:shader,execution,memory_model,atomicity:atomicity:*": { "subcaseMS": 77.201 },
   "webgpu:shader,execution,memory_model,barrier:workgroup_barrier_load_store:*": { "subcaseMS": 65.850 },
@@ -1682,7 +1692,7 @@
   "webgpu:shader,execution,memory_model,coherence:corw2:*": { "subcaseMS": 244.384 },
   "webgpu:shader,execution,memory_model,coherence:cowr:*": { "subcaseMS": 250.484 },
   "webgpu:shader,execution,memory_model,coherence:coww:*": { "subcaseMS": 245.850 },
-  "webgpu:shader,execution,memory_model,texture_intra_invocation_coherence:texture_intra_invocation_coherence:*": { "subcaseMS": 0 },
+  "webgpu:shader,execution,memory_model,texture_intra_invocation_coherence:texture_intra_invocation_coherence:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,memory_model,weak:2_plus_2_write:*": { "subcaseMS": 185.150 },
   "webgpu:shader,execution,memory_model,weak:load_buffer:*": { "subcaseMS": 184.900 },
   "webgpu:shader,execution,memory_model,weak:message_passing:*": { "subcaseMS": 196.550 },
@@ -1701,15 +1711,15 @@
   "webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:*": { "subcaseMS": 6.487 },
   "webgpu:shader,execution,shader_io,compute_builtins:inputs:*": { "subcaseMS": 19.342 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,front_facing:*": { "subcaseMS": 1.001 },
-  "webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage,centroid:*": { "subcaseMS": 1.001 },
+  "webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,position:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,sample_index:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,sample_mask:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,shared_structs:shared_between_stages:*": { "subcaseMS": 9.601 },
   "webgpu:shader,execution,shader_io,shared_structs:shared_with_buffer:*": { "subcaseMS": 20.701 },
   "webgpu:shader,execution,shader_io,shared_structs:shared_with_non_entry_point_function:*": { "subcaseMS": 6.801 },
-  "webgpu:shader,execution,shader_io,workgroup_size:workgroup_size:*": { "subcaseMS": 0 },
+  "webgpu:shader,execution,shader_io,workgroup_size:workgroup_size:*": { "subcaseMS": 0.000 },
   "webgpu:shader,execution,shadow:builtin:*": { "subcaseMS": 4.700 },
   "webgpu:shader,execution,shadow:declaration:*": { "subcaseMS": 9.700 },
   "webgpu:shader,execution,shadow:for_loop:*": { "subcaseMS": 17.201 },
@@ -1758,17 +1768,17 @@
   "webgpu:shader,validation,decl,ptr_spelling:ptr_bad_store_type:*": { "subcaseMS": 0.967 },
   "webgpu:shader,validation,decl,ptr_spelling:ptr_handle_space_invalid:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,ptr_spelling:ptr_not_instantiable:*": { "subcaseMS": 1.310 },
-  "webgpu:shader,validation,decl,var:binding_point_on_resources:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:binding_point_on_non_resources:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:binding_point_on_function_var:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:binding_collisions:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:binding_collision_unused_helper:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:module_scope_types:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:function_scope_types:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_collisions:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_point_on_function_var:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_point_on_non_resources:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_point_on_resources:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:function_addrspace_at_module_scope:*": { "subcaseMS": 1.000 },
-  "webgpu:shader,validation,decl,var:module_scope_initializers:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:function_scope_types:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:handle_initializer:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:initializer_kind:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:module_scope_initializers:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:module_scope_types:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var_access_mode:explicit_access_mode:*": { "subcaseMS": 1.373 },
   "webgpu:shader,validation,decl,var_access_mode:implicit_access_mode:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var_access_mode:read_access:*": { "subcaseMS": 1.177 },
@@ -1783,6 +1793,9 @@
   "webgpu:shader,validation,expression,call,builtin,acos:values:*": { "subcaseMS": 0.342 },
   "webgpu:shader,validation,expression,call,builtin,acosh:integer_argument:*": { "subcaseMS": 1.234 },
   "webgpu:shader,validation,expression,call,builtin,acosh:values:*": { "subcaseMS": 0.217 },
+  "webgpu:shader,validation,expression,call,builtin,arrayLength:access_mode:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,expression,call,builtin,arrayLength:bool_type:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,expression,call,builtin,arrayLength:type:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,expression,call,builtin,asin:integer_argument:*": { "subcaseMS": 0.878 },
   "webgpu:shader,validation,expression,call,builtin,asin:values:*": { "subcaseMS": 0.359 },
   "webgpu:shader,validation,expression,call,builtin,asinh:integer_argument:*": { "subcaseMS": 1.267 },
@@ -1795,8 +1808,8 @@
   "webgpu:shader,validation,expression,call,builtin,atanh:integer_argument:*": { "subcaseMS": 0.912 },
   "webgpu:shader,validation,expression,call,builtin,atanh:values:*": { "subcaseMS": 0.231 },
   "webgpu:shader,validation,expression,call,builtin,atomics:stage:*": { "subcaseMS": 1.346 },
-  "webgpu:shader,validation,expression,call,builtin,barriers:only_in_compute:*": { "subcaseMS": 1.500 },
   "webgpu:shader,validation,expression,call,builtin,barriers:no_return_value:*": { "subcaseMS": 1.500 },
+  "webgpu:shader,validation,expression,call,builtin,barriers:only_in_compute:*": { "subcaseMS": 1.500 },
   "webgpu:shader,validation,expression,call,builtin,bitcast:bad_const_to_f16:*": { "subcaseMS": 0.753 },
   "webgpu:shader,validation,expression,call,builtin,bitcast:bad_const_to_f32:*": { "subcaseMS": 0.844 },
   "webgpu:shader,validation,expression,call,builtin,bitcast:bad_to_f16:*": { "subcaseMS": 8.518 },
@@ -1881,26 +1894,29 @@
   "webgpu:shader,validation,expression,call,builtin,unpack4xU8:must_use:*": { "subcaseMS": 32.800 },
   "webgpu:shader,validation,expression,call,builtin,unpack4xU8:supported:*": { "subcaseMS": 98.501 },
   "webgpu:shader,validation,expression,call,builtin,unpack4xU8:unsupported:*": { "subcaseMS": 346.801 },
-  "webgpu:shader,validation,extension,pointer_composite_access:deref:*": { "subcaseMS": 0.0 },
-  "webgpu:shader,validation,extension,pointer_composite_access:pointer:*": { "subcaseMS": 0.0 },
+  "webgpu:shader,validation,expression,unary,address_of_and_indirection:basic:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,expression,unary,address_of_and_indirection:composite:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,expression,unary,address_of_and_indirection:invalid:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,extension,pointer_composite_access:deref:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,extension,pointer_composite_access:pointer:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,functions,alias_analysis:aliasing_inside_function:*": { "subcaseMS": 1.200 },
   "webgpu:shader,validation,functions,alias_analysis:member_accessors:*": { "subcaseMS": 1.656 },
-  "webgpu:shader,validation,functions,alias_analysis:one_atomic_pointer_one_module_scope:*": { "subcaseMS": 0 },
+  "webgpu:shader,validation,functions,alias_analysis:one_atomic_pointer_one_module_scope:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,functions,alias_analysis:one_pointer_one_module_scope:*": { "subcaseMS": 1.598 },
   "webgpu:shader,validation,functions,alias_analysis:same_pointer_read_and_write:*": { "subcaseMS": 1.301 },
   "webgpu:shader,validation,functions,alias_analysis:subcalls:*": { "subcaseMS": 1.673 },
-  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers_to_array_elements:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers_to_struct_members:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_array_elements_indirect:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_array_elements:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_struct_members_indirect:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_struct_members:*": { "subcaseMS": 0 },
+  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers_to_array_elements:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers_to_struct_members:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,functions,alias_analysis:two_pointers:*": { "subcaseMS": 1.537 },
-  "webgpu:shader,validation,functions,alias_analysis:workgroup_uniform_load:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_1_param:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_2_params:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_3_params:*": { "subcaseMS": 0 },
+  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_array_elements:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_array_elements_indirect:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_struct_members:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:two_pointers_to_struct_members_indirect:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,alias_analysis:workgroup_uniform_load:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_1_param:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_2_params:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,functions,restrictions:call_arg_types_match_3_params:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,functions,restrictions:entry_point_call_target:*": { "subcaseMS": 1.734 },
   "webgpu:shader,validation,functions,restrictions:function_parameter_matching:*": { "subcaseMS": 1.953 },
   "webgpu:shader,validation,functions,restrictions:function_parameter_types:*": { "subcaseMS": 1.520 },
@@ -1921,6 +1937,8 @@
   "webgpu:shader,validation,parse,blankspace:bom:*": { "subcaseMS": 1.101 },
   "webgpu:shader,validation,parse,blankspace:null_characters:*": { "subcaseMS": 3.217 },
   "webgpu:shader,validation,parse,break:placement:*": { "subcaseMS": 1.254 },
+  "webgpu:shader,validation,parse,break_if:non_bool_param:*": { "subcaseMS": 0.000 },
+  "webgpu:shader,validation,parse,break_if:placement:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,parse,builtin:parse:*": { "subcaseMS": 3.277 },
   "webgpu:shader,validation,parse,builtin:placement:*": { "subcaseMS": 1.267 },
   "webgpu:shader,validation,parse,comments:comments:*": { "subcaseMS": 1.000 },
@@ -1930,6 +1948,7 @@
   "webgpu:shader,validation,parse,compound:parse:*": { "subcaseMS": 4.315 },
   "webgpu:shader,validation,parse,const:placement:*": { "subcaseMS": 1.167 },
   "webgpu:shader,validation,parse,const_assert:parse:*": { "subcaseMS": 1.400 },
+  "webgpu:shader,validation,parse,continuing:placement:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,parse,diagnostic:after_other_directives:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,diagnostic:conflicting_attribute_different_location:*": { "subcaseMS": 2.257 },
   "webgpu:shader,validation,parse,diagnostic:conflicting_attribute_same_location:*": { "subcaseMS": 1.400 },
@@ -1963,7 +1982,7 @@
   "webgpu:shader,validation,parse,must_use:builtin_no_must_use:*": { "subcaseMS": 1.206 },
   "webgpu:shader,validation,parse,must_use:call:*": { "subcaseMS": 1.275 },
   "webgpu:shader,validation,parse,must_use:declaration:*": { "subcaseMS": 1.523 },
-  "webgpu:shader,validation,parse,must_use:ignore_result_of_non_must_use_that_returns_call_of_must_use:*": { "subcaseMS": 0.0 },
+  "webgpu:shader,validation,parse,must_use:ignore_result_of_non_must_use_that_returns_call_of_must_use:*": { "subcaseMS": 0.000 },
   "webgpu:shader,validation,parse,pipeline_stage:compute_parsing:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,parse,pipeline_stage:extra_on_compute_function:*": { "subcaseMS": 2.651 },
   "webgpu:shader,validation,parse,pipeline_stage:extra_on_fragment_function:*": { "subcaseMS": 1.001 },
@@ -2060,14 +2079,14 @@
   "webgpu:shader,validation,shader_io,locations:duplicates:*": { "subcaseMS": 1.906 },
   "webgpu:shader,validation,shader_io,locations:location_fp16:*": { "subcaseMS": 0.501 },
   "webgpu:shader,validation,shader_io,locations:nesting:*": { "subcaseMS": 0.967 },
+  "webgpu:shader,validation,shader_io,locations:out_of_order:*": { "subcaseMS": 1.296 },
   "webgpu:shader,validation,shader_io,locations:stage_inout:*": { "subcaseMS": 1.850 },
   "webgpu:shader,validation,shader_io,locations:type:*": { "subcaseMS": 1.332 },
   "webgpu:shader,validation,shader_io,locations:validation:*": { "subcaseMS": 1.296 },
-  "webgpu:shader,validation,shader_io,locations:out_of_order:*": { "subcaseMS": 1.296 },
   "webgpu:shader,validation,shader_io,size:size:*": { "subcaseMS": 1.218 },
+  "webgpu:shader,validation,shader_io,size:size_creation_fixed_footprint:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,shader_io,size:size_fp16:*": { "subcaseMS": 1.500 },
   "webgpu:shader,validation,shader_io,size:size_non_struct:*": { "subcaseMS": 0.929 },
-  "webgpu:shader,validation,shader_io,size:size_creation_fixed_footprint:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size:*": { "subcaseMS": 1.227 },
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size_const:*": { "subcaseMS": 3.400 },
   "webgpu:shader,validation,shader_io,workgroup_size:workgroup_size_fp16:*": { "subcaseMS": 0.700 },
@@ -2151,25 +2170,7 @@
   "webgpu:web_platform,external_texture,video:importExternalTexture,compute:*": { "subcaseMS": 36.270 },
   "webgpu:web_platform,external_texture,video:importExternalTexture,sample:*": { "subcaseMS": 34.968 },
   "webgpu:web_platform,external_texture,video:importExternalTexture,sampleWithVideoFrameWithVisibleRectParam:*": { "subcaseMS": 29.160 },
-  "webgpu:web_platform,worker,worker:worker:*": { "subcaseMS": 245.901 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_f32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_i32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_u32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:af_to_vec2f16:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:vec2af_to_vec4f16:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_i32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_u32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_f32:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:ai_to_vec2f16:*": { "subcaseMS": 0 },
-  "webgpu:shader,execution,expression,call,builtin,bitcast:vec2ai_to_vec4f16:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,parse,continuing:placement:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,parse,break_if:placement:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,parse,break_if:non_bool_param:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,call,builtin,arrayLength:bool_type:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,call,builtin,arrayLength:type:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,call,builtin,arrayLength:access_mode:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,unary,address_of_and_indirection:basic:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,unary,address_of_and_indirection:composite:*": { "subcaseMS": 0 },
-  "webgpu:shader,validation,expression,unary,address_of_and_indirection:invalid:*": { "subcaseMS": 0 },
+  "webgpu:web_platform,worker,worker:dedicated_worker:*": { "subcaseMS": 245.901 },
+  "webgpu:web_platform,worker,worker:shared_worker:*": { "subcaseMS": 26.801 },
   "_end": ""
 }
diff --git a/src/webgpu/web_platform/worker/worker.spec.ts b/src/webgpu/web_platform/worker/worker.spec.ts
index 67f9f693bee0..0f34a58ddaa5 100644
--- a/src/webgpu/web_platform/worker/worker.spec.ts
+++ b/src/webgpu/web_platform/worker/worker.spec.ts
@@ -1,9 +1,9 @@
 export const description = `
-Tests WebGPU is available in a worker.
+Tests WebGPU is available in a dedicated worker and a shared worker.
 
-Note: The CTS test can be run in a worker by passing in worker=1 as
-a query parameter. This test is specifically to check that WebGPU
-is available in a worker.
+Note: The CTS test can be run respectively in a dedicated worker and a shared worker by
+passing in worker=dedicated and worker=shared as a query parameter. These tests
+are specifically to check that WebGPU is available in a dedicated worker and a shared worker.
 `;
 
 import { Fixture } from '../../../common/framework/fixture.js';
@@ -16,8 +16,8 @@ function isNode(): boolean {
   return typeof process !== 'undefined' && process?.versions?.node !== undefined;
 }
 
-g.test('worker')
-  .desc(`test WebGPU is available in DedicatedWorkers and check for basic functionality`)
+g.test('dedicated_worker')
+  .desc(`test WebGPU is available in dedicated workers and check for basic functionality`)
   .fn(async t => {
     if (isNode()) {
       t.skip('node does not support 100% compatible workers');
@@ -29,7 +29,25 @@ g.test('worker')
     // otherwise typescript tries to parse the file which again, fails.
     // worker_launcher.js is excluded in node.tsconfig.json.
     const url = './worker_launcher.js';
-    const { launchWorker } = await import(url);
-    const result = await launchWorker();
+    const { launchDedicatedWorker } = await import(url);
+    const result = await launchDedicatedWorker();
+    assert(result.error === undefined, `should be no error from worker but was: ${result.error}`);
+  });
+
+g.test('shared_worker')
+  .desc(`test WebGPU is available in shared workers and check for basic functionality`)
+  .fn(async t => {
+    if (isNode()) {
+      t.skip('node does not support 100% compatible workers');
+      return;
+    }
+    // Note: we load worker_launcher dynamically because ts-node support
+    // is using commonjs which doesn't support import.meta. Further,
+    // we need to put the url in a string add pass the string to import
+    // otherwise typescript tries to parse the file which again, fails.
+    // worker_launcher.js is excluded in node.tsconfig.json.
+    const url = './worker_launcher.js';
+    const { launchSharedWorker } = await import(url);
+    const result = await launchSharedWorker();
     assert(result.error === undefined, `should be no error from worker but was: ${result.error}`);
   });
diff --git a/src/webgpu/web_platform/worker/worker.ts b/src/webgpu/web_platform/worker/worker.ts
index a3cf8064e26a..f3c907a4118a 100644
--- a/src/webgpu/web_platform/worker/worker.ts
+++ b/src/webgpu/web_platform/worker/worker.ts
@@ -1,6 +1,10 @@
 import { getGPU, setDefaultRequestAdapterOptions } from '../../../common/util/navigator_gpu.js';
 import { assert, objectEquals, iterRange } from '../../../common/util/util.js';
 
+// Should be WorkerGlobalScope, but importing lib "webworker" conflicts with lib "dom".
+/* eslint-disable-next-line @typescript-eslint/no-explicit-any */
+declare const self: any;
+
 async function basicTest() {
   const adapter = await getGPU(null).requestAdapter();
   assert(adapter !== null, 'Failed to get adapter.');
@@ -68,7 +72,7 @@ async function basicTest() {
   device.destroy();
 }
 
-self.onmessage = async (ev: MessageEvent) => {
+async function reportTestResults(this: MessagePort | Worker, ev: MessageEvent) {
   const defaultRequestAdapterOptions: GPURequestAdapterOptions =
     ev.data.defaultRequestAdapterOptions;
   setDefaultRequestAdapterOptions(defaultRequestAdapterOptions);
@@ -79,5 +83,17 @@ self.onmessage = async (ev: MessageEvent) => {
   } catch (err: unknown) {
     error = (err as Error).toString();
   }
-  self.postMessage({ error });
+  this.postMessage({ error });
+}
+
+self.onmessage = (ev: MessageEvent) => {
+  void reportTestResults.call(self, ev);
+};
+
+self.onconnect = (event: MessageEvent) => {
+  const port = event.ports[0];
+
+  port.onmessage = (ev: MessageEvent) => {
+    void reportTestResults.call(port, ev);
+  };
 };
diff --git a/src/webgpu/web_platform/worker/worker_launcher.ts b/src/webgpu/web_platform/worker/worker_launcher.ts
index 72059eb99fda..0487e4ad38b4 100644
--- a/src/webgpu/web_platform/worker/worker_launcher.ts
+++ b/src/webgpu/web_platform/worker/worker_launcher.ts
@@ -4,7 +4,7 @@ export type TestResult = {
   error: String | undefined;
 };
 
-export async function launchWorker() {
+export async function launchDedicatedWorker() {
   const selfPath = import.meta.url;
   const selfPathDir = selfPath.substring(0, selfPath.lastIndexOf('/'));
   const workerPath = selfPathDir + '/worker.js';
@@ -16,3 +16,20 @@ export async function launchWorker() {
   worker.postMessage({ defaultRequestAdapterOptions: getDefaultRequestAdapterOptions() });
   return await promise;
 }
+
+export async function launchSharedWorker() {
+  const selfPath = import.meta.url;
+  const selfPathDir = selfPath.substring(0, selfPath.lastIndexOf('/'));
+  const workerPath = selfPathDir + '/worker.js';
+  const worker = new SharedWorker(workerPath, { type: 'module' });
+
+  const port = worker.port;
+  const promise = new Promise<TestResult>(resolve => {
+    port.addEventListener('message', ev => resolve(ev.data as TestResult), { once: true });
+  });
+  port.start();
+  port.postMessage({
+    defaultRequestAdapterOptions: getDefaultRequestAdapterOptions(),
+  });
+  return await promise;
+}