diff --git a/docs/intro/developing.md b/docs/intro/developing.md
index 23da946b1e0c..c706e0c1c1b8 100644
--- a/docs/intro/developing.md
+++ b/docs/intro/developing.md
@@ -56,7 +56,7 @@ 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=dedicated` runs the tests on a dedicated worker instead of the main thread.
+- `worker=dedicated` (or `worker` or `worker=1`) 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.
 - `worker=service` runs the tests on a service worker instead of the main thread.
 - `power_preference=low-power` runs most tests passing `powerPreference: low-power` to `requestAdapter`
diff --git a/src/common/internal/query/query.ts b/src/common/internal/query/query.ts
index ad7cf246e9d9..676ac46d3888 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 { optionString } from '../../runtime/helper/options.js';
+import { optionWorkerMode } from '../../runtime/helper/options.js';
 import { assert, unreachable } from '../../util/util.js';
 import { Expectation } from '../logging/result.js';
 
@@ -193,7 +193,7 @@ Expectation should be of the form path/to/cts.https.html?debug=0&q=suite:test_pa
       );
 
       const params = expectationURL.searchParams;
-      if (optionString('worker', params) !== optionString('worker', wptURL.searchParams)) {
+      if (optionWorkerMode('worker', params) !== optionWorkerMode('worker', wptURL.searchParams)) {
         continue;
       }
 
diff --git a/src/common/runtime/helper/options.ts b/src/common/runtime/helper/options.ts
index 67fd00372d39..38000341d945 100644
--- a/src/common/runtime/helper/options.ts
+++ b/src/common/runtime/helper/options.ts
@@ -1,3 +1,5 @@
+import { unreachable } from '../../util/util.js';
+
 let windowURL: URL | undefined = undefined;
 function getWindowURL() {
   if (windowURL === undefined) {
@@ -6,6 +8,7 @@ function getWindowURL() {
   return windowURL;
 }
 
+/** Parse a runner option that is always boolean-typed. False if missing or '0'. */
 export function optionEnabled(
   opt: string,
   searchParams: URLSearchParams = getWindowURL().searchParams
@@ -14,6 +17,7 @@ export function optionEnabled(
   return val !== null && val !== '0';
 }
 
+/** Parse a runner option that is always string-typed. If the option is missing, returns `''`. */
 export function optionString(
   opt: string,
   searchParams: URLSearchParams = getWindowURL().searchParams
@@ -21,20 +25,40 @@ export function optionString(
   return searchParams.get(opt) || '';
 }
 
+/** Runtime modes for whether to run tests in a worker. '0' means no worker. */
+type WorkerMode = '0' | 'dedicated' | 'service' | 'shared';
+/** Parse a runner option for different worker modes (as in `?worker=shared`). */
+export function optionWorkerMode(
+  opt: string,
+  searchParams: URLSearchParams = getWindowURL().searchParams
+): WorkerMode {
+  const value = searchParams.get(opt);
+  if (value === null || value === '0') {
+    return '0';
+  } else if (value === 'service') {
+    return 'service';
+  } else if (value === 'shared') {
+    return 'shared';
+  } else if (value === '' || value === '1' || value === 'dedicated') {
+    return 'dedicated';
+  }
+  unreachable('invalid worker= option value');
+}
+
 /**
  * The possible options for the tests.
  */
 export interface CTSOptions {
-  worker?: 'dedicated' | 'shared' | 'service' | '';
+  worker: WorkerMode;
   debug: boolean;
   compatibility: boolean;
   forceFallbackAdapter: boolean;
   unrollConstEvalLoops: boolean;
-  powerPreference?: GPUPowerPreference | '';
+  powerPreference: GPUPowerPreference | '';
 }
 
 export const kDefaultCTSOptions: CTSOptions = {
-  worker: '',
+  worker: '0',
   debug: true,
   compatibility: false,
   forceFallbackAdapter: false,
@@ -63,9 +87,9 @@ export type OptionsInfos<Type> = Record<keyof Type, OptionInfo>;
 export const kCTSOptionsInfo: OptionsInfos<CTSOptions> = {
   worker: {
     description: 'run in a worker',
-    parser: optionString,
+    parser: optionWorkerMode,
     selectValueDescriptions: [
-      { value: '', description: 'no worker' },
+      { value: '0', description: 'no worker' },
       { value: 'dedicated', description: 'dedicated worker' },
       { value: 'shared', description: 'shared worker' },
       { value: 'service', description: 'service worker' },
diff --git a/src/common/runtime/standalone.ts b/src/common/runtime/standalone.ts
index 62090ee5dc92..385de9937e09 100644
--- a/src/common/runtime/standalone.ts
+++ b/src/common/runtime/standalone.ts
@@ -550,7 +550,7 @@ function keyValueToPairs([k, v]: [string, ParamValue]): [string, string][] {
  */
 function prepareParams(params: Record<string, ParamValue>): string {
   const pairsArrays = Object.entries(params)
-    .filter(([, v]) => !!v)
+    .filter(([, v]) => !!v && v !== '0')
     .map(keyValueToPairs);
   const pairs = pairsArrays.flat();
   return new URLSearchParams(pairs).toString();
diff --git a/src/common/runtime/wpt.ts b/src/common/runtime/wpt.ts
index 56ff9649e045..79ed1b5924dd 100644
--- a/src/common/runtime/wpt.ts
+++ b/src/common/runtime/wpt.ts
@@ -8,7 +8,7 @@ import { parseQuery } from '../internal/query/parseQuery.js';
 import { parseExpectationsForTestQuery, relativeQueryString } from '../internal/query/query.js';
 import { assert } from '../util/util.js';
 
-import { optionEnabled, optionString } from './helper/options.js';
+import { optionEnabled, optionWorkerMode } from './helper/options.js';
 import { TestDedicatedWorker, TestServiceWorker, TestSharedWorker } from './helper/test_worker.js';
 
 // testharness.js API (https://web-platform-tests.org/writing-tests/testharness-api.html)
@@ -31,7 +31,7 @@ setup({
 });
 
 void (async () => {
-  const workerString = optionString('worker');
+  const workerString = optionWorkerMode('worker');
   const dedicatedWorker = workerString === 'dedicated' ? new TestDedicatedWorker() : undefined;
   const sharedWorker = workerString === 'shared' ? new TestSharedWorker() : undefined;
   const serviceWorker = workerString === 'service' ? new TestServiceWorker() : undefined;
diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json
index 1dfb7379313f..6a616d571a9e 100644
--- a/src/resources/cache/hashes.json
+++ b/src/resources/cache/hashes.json
@@ -1,110 +1,110 @@
 {
-  "webgpu/shader/execution/binary/af_addition.bin": "f79f0374",
-  "webgpu/shader/execution/binary/af_logical.bin": "9ab41597",
-  "webgpu/shader/execution/binary/af_division.bin": "770fc018",
-  "webgpu/shader/execution/binary/af_matrix_addition.bin": "a401d81a",
-  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "a9e51912",
-  "webgpu/shader/execution/binary/af_multiplication.bin": "8f7b5f84",
-  "webgpu/shader/execution/binary/af_remainder.bin": "133a1dd7",
-  "webgpu/shader/execution/binary/af_subtraction.bin": "9148f3ce",
-  "webgpu/shader/execution/binary/f16_addition.bin": "d58a9371",
-  "webgpu/shader/execution/binary/f16_logical.bin": "80f9e754",
-  "webgpu/shader/execution/binary/f16_division.bin": "5d2bf113",
-  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "531954a6",
-  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "b871cd18",
-  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "3603d1db",
-  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "c97537aa",
-  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "273b76c6",
-  "webgpu/shader/execution/binary/f16_multiplication.bin": "d98cd937",
-  "webgpu/shader/execution/binary/f16_remainder.bin": "6e5227f4",
-  "webgpu/shader/execution/binary/f16_subtraction.bin": "d243bbad",
-  "webgpu/shader/execution/binary/f32_addition.bin": "c288aef6",
-  "webgpu/shader/execution/binary/f32_logical.bin": "47467130",
-  "webgpu/shader/execution/binary/f32_division.bin": "c9c76117",
-  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "88eedbfa",
-  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "9f76f2ac",
-  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "1d667df9",
-  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "aa253cf1",
-  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "98a63b3a",
-  "webgpu/shader/execution/binary/f32_multiplication.bin": "dce8dd5b",
-  "webgpu/shader/execution/binary/f32_remainder.bin": "d3bdb9e3",
-  "webgpu/shader/execution/binary/f32_subtraction.bin": "8a48e67a",
-  "webgpu/shader/execution/binary/i32_arithmetic.bin": "d345f9c7",
-  "webgpu/shader/execution/binary/i32_comparison.bin": "6d4ae7e0",
-  "webgpu/shader/execution/binary/u32_arithmetic.bin": "163f8a38",
-  "webgpu/shader/execution/binary/u32_comparison.bin": "f6c1497b",
-  "webgpu/shader/execution/abs.bin": "e21a06df",
-  "webgpu/shader/execution/acos.bin": "2a92d3d5",
-  "webgpu/shader/execution/acosh.bin": "c2cc2e23",
-  "webgpu/shader/execution/asin.bin": "2027460e",
-  "webgpu/shader/execution/asinh.bin": "21bac67b",
-  "webgpu/shader/execution/atan.bin": "6422056c",
-  "webgpu/shader/execution/atan2.bin": "d8ade832",
-  "webgpu/shader/execution/atanh.bin": "de1dd54c",
-  "webgpu/shader/execution/bitcast.bin": "fdd0874",
-  "webgpu/shader/execution/ceil.bin": "6596b2e2",
-  "webgpu/shader/execution/clamp.bin": "9a92d1b7",
-  "webgpu/shader/execution/cos.bin": "647f5fda",
-  "webgpu/shader/execution/cosh.bin": "5f847611",
-  "webgpu/shader/execution/cross.bin": "117f2e0f",
-  "webgpu/shader/execution/degrees.bin": "cadaa756",
-  "webgpu/shader/execution/determinant.bin": "f2df9222",
-  "webgpu/shader/execution/distance.bin": "d6a161d1",
-  "webgpu/shader/execution/dot.bin": "c788d592",
-  "webgpu/shader/execution/exp.bin": "a76c67ca",
-  "webgpu/shader/execution/exp2.bin": "a5b443c5",
-  "webgpu/shader/execution/faceForward.bin": "75f1f377",
-  "webgpu/shader/execution/floor.bin": "8aea5f9c",
-  "webgpu/shader/execution/fma.bin": "8682909f",
-  "webgpu/shader/execution/fract.bin": "ff3283f0",
-  "webgpu/shader/execution/frexp.bin": "9639418e",
-  "webgpu/shader/execution/inverseSqrt.bin": "9de3fa15",
-  "webgpu/shader/execution/ldexp.bin": "43fde30e",
-  "webgpu/shader/execution/length.bin": "24b87c1a",
-  "webgpu/shader/execution/log.bin": "c8447a16",
-  "webgpu/shader/execution/log2.bin": "bcefcef6",
-  "webgpu/shader/execution/max.bin": "bbd522ab",
-  "webgpu/shader/execution/min.bin": "4687f034",
-  "webgpu/shader/execution/mix.bin": "e79fc2e8",
-  "webgpu/shader/execution/modf.bin": "feedecc3",
-  "webgpu/shader/execution/normalize.bin": "78e7c0",
-  "webgpu/shader/execution/pack2x16float.bin": "e472a927",
-  "webgpu/shader/execution/pow.bin": "38d0ff7",
-  "webgpu/shader/execution/quantizeToF16.bin": "5496ab3",
-  "webgpu/shader/execution/radians.bin": "db736290",
-  "webgpu/shader/execution/reflect.bin": "b08379ba",
-  "webgpu/shader/execution/refract.bin": "d360cf82",
-  "webgpu/shader/execution/round.bin": "3d2c318e",
-  "webgpu/shader/execution/saturate.bin": "33892d9c",
-  "webgpu/shader/execution/sign.bin": "8a54e6ec",
-  "webgpu/shader/execution/sin.bin": "f0cd3a32",
-  "webgpu/shader/execution/sinh.bin": "4d1b0134",
-  "webgpu/shader/execution/smoothstep.bin": "a3b004be",
-  "webgpu/shader/execution/sqrt.bin": "dd8a5970",
-  "webgpu/shader/execution/step.bin": "eec1579d",
-  "webgpu/shader/execution/tan.bin": "612d43b",
-  "webgpu/shader/execution/tanh.bin": "6f95242c",
-  "webgpu/shader/execution/transpose.bin": "b1a9a7f7",
-  "webgpu/shader/execution/trunc.bin": "cdf5939e",
-  "webgpu/shader/execution/unpack2x16float.bin": "b992957a",
-  "webgpu/shader/execution/unpack2x16snorm.bin": "394e0942",
-  "webgpu/shader/execution/unpack2x16unorm.bin": "78f0ef3c",
-  "webgpu/shader/execution/unpack4x8snorm.bin": "926f3f3e",
-  "webgpu/shader/execution/unpack4x8unorm.bin": "3640a996",
-  "webgpu/shader/execution/unary/af_arithmetic.bin": "bf05a1b1",
-  "webgpu/shader/execution/unary/af_assignment.bin": "d711e244",
-  "webgpu/shader/execution/unary/bool_conversion.bin": "d0c1e5a3",
-  "webgpu/shader/execution/unary/f16_arithmetic.bin": "6095e89f",
-  "webgpu/shader/execution/unary/f16_conversion.bin": "76546200",
-  "webgpu/shader/execution/unary/f32_arithmetic.bin": "a16e7cf4",
-  "webgpu/shader/execution/unary/f32_conversion.bin": "62604b96",
-  "webgpu/shader/execution/unary/i32_arithmetic.bin": "a8649cbb",
-  "webgpu/shader/execution/unary/i32_conversion.bin": "e5157a69",
-  "webgpu/shader/execution/unary/u32_conversion.bin": "d07d0c20",
-  "webgpu/shader/execution/unary/ai_assignment.bin": "f62c765c",
-  "webgpu/shader/execution/binary/ai_arithmetic.bin": "43501242",
-  "webgpu/shader/execution/unary/ai_arithmetic.bin": "8e448c53",
-  "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "c57c159e",
-  "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "d21017f3",
-  "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "ddd2f83a"
+  "webgpu/shader/execution/binary/af_addition.bin": "24160909",
+  "webgpu/shader/execution/binary/af_logical.bin": "ffb5a83f",
+  "webgpu/shader/execution/binary/af_division.bin": "c230ac78",
+  "webgpu/shader/execution/binary/af_matrix_addition.bin": "9079a042",
+  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "6b55102f",
+  "webgpu/shader/execution/binary/af_multiplication.bin": "4fc3b0d6",
+  "webgpu/shader/execution/binary/af_remainder.bin": "366caec6",
+  "webgpu/shader/execution/binary/af_subtraction.bin": "49a16db4",
+  "webgpu/shader/execution/binary/f16_addition.bin": "3fb1ee09",
+  "webgpu/shader/execution/binary/f16_logical.bin": "65cdc6f",
+  "webgpu/shader/execution/binary/f16_division.bin": "1b5bd44a",
+  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "fdea291a",
+  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "e727482e",
+  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "24d70bdd",
+  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "7acd3c3b",
+  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "9e01e0cf",
+  "webgpu/shader/execution/binary/f16_multiplication.bin": "c24b705c",
+  "webgpu/shader/execution/binary/f16_remainder.bin": "2764c7e1",
+  "webgpu/shader/execution/binary/f16_subtraction.bin": "d4014a38",
+  "webgpu/shader/execution/binary/f32_addition.bin": "b6707259",
+  "webgpu/shader/execution/binary/f32_logical.bin": "d40c1c0",
+  "webgpu/shader/execution/binary/f32_division.bin": "e19e30a6",
+  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "6a13d6d6",
+  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "4c5cb0a2",
+  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "6b00d17f",
+  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "319d3ae1",
+  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "1c4893ca",
+  "webgpu/shader/execution/binary/f32_multiplication.bin": "f7dbbc8",
+  "webgpu/shader/execution/binary/f32_remainder.bin": "fb5dd3fe",
+  "webgpu/shader/execution/binary/f32_subtraction.bin": "aeaca568",
+  "webgpu/shader/execution/binary/i32_arithmetic.bin": "90a3514a",
+  "webgpu/shader/execution/binary/i32_comparison.bin": "840b09e6",
+  "webgpu/shader/execution/binary/u32_arithmetic.bin": "e3f04e97",
+  "webgpu/shader/execution/binary/u32_comparison.bin": "dc1398d3",
+  "webgpu/shader/execution/abs.bin": "2b663377",
+  "webgpu/shader/execution/acos.bin": "a888c67f",
+  "webgpu/shader/execution/acosh.bin": "5383658b",
+  "webgpu/shader/execution/asin.bin": "a580a322",
+  "webgpu/shader/execution/asinh.bin": "31388535",
+  "webgpu/shader/execution/atan.bin": "c6f4c771",
+  "webgpu/shader/execution/atan2.bin": "3b9a37a1",
+  "webgpu/shader/execution/atanh.bin": "7295a313",
+  "webgpu/shader/execution/bitcast.bin": "73e0bea6",
+  "webgpu/shader/execution/ceil.bin": "24bd6be9",
+  "webgpu/shader/execution/clamp.bin": "8ed55492",
+  "webgpu/shader/execution/cos.bin": "4d36fa0",
+  "webgpu/shader/execution/cosh.bin": "a7da3a3",
+  "webgpu/shader/execution/cross.bin": "6ce2660b",
+  "webgpu/shader/execution/degrees.bin": "53a0849",
+  "webgpu/shader/execution/determinant.bin": "873e78af",
+  "webgpu/shader/execution/distance.bin": "bfce09a0",
+  "webgpu/shader/execution/dot.bin": "566b6e55",
+  "webgpu/shader/execution/exp.bin": "d8142d49",
+  "webgpu/shader/execution/exp2.bin": "52a403d1",
+  "webgpu/shader/execution/faceForward.bin": "f90dacbe",
+  "webgpu/shader/execution/floor.bin": "c7465b11",
+  "webgpu/shader/execution/fma.bin": "db3497d4",
+  "webgpu/shader/execution/fract.bin": "e65c5bd1",
+  "webgpu/shader/execution/frexp.bin": "5ce112f4",
+  "webgpu/shader/execution/inverseSqrt.bin": "84ea0a57",
+  "webgpu/shader/execution/ldexp.bin": "9d11f120",
+  "webgpu/shader/execution/length.bin": "dd759736",
+  "webgpu/shader/execution/log.bin": "5f77a59d",
+  "webgpu/shader/execution/log2.bin": "57fff45e",
+  "webgpu/shader/execution/max.bin": "c7cdd54f",
+  "webgpu/shader/execution/min.bin": "422be325",
+  "webgpu/shader/execution/mix.bin": "957e7d92",
+  "webgpu/shader/execution/modf.bin": "946d1f4f",
+  "webgpu/shader/execution/normalize.bin": "2ac02e80",
+  "webgpu/shader/execution/pack2x16float.bin": "d7ef3cf5",
+  "webgpu/shader/execution/pow.bin": "cd492166",
+  "webgpu/shader/execution/quantizeToF16.bin": "58bac06c",
+  "webgpu/shader/execution/radians.bin": "6ac08f50",
+  "webgpu/shader/execution/reflect.bin": "3c260554",
+  "webgpu/shader/execution/refract.bin": "7e952d7c",
+  "webgpu/shader/execution/round.bin": "f6b9bda1",
+  "webgpu/shader/execution/saturate.bin": "1c22f301",
+  "webgpu/shader/execution/sign.bin": "3ce55105",
+  "webgpu/shader/execution/sin.bin": "d48e7f2a",
+  "webgpu/shader/execution/sinh.bin": "f227c1d1",
+  "webgpu/shader/execution/smoothstep.bin": "6bdb4309",
+  "webgpu/shader/execution/sqrt.bin": "cd576d27",
+  "webgpu/shader/execution/step.bin": "dd584686",
+  "webgpu/shader/execution/tan.bin": "5ae50f61",
+  "webgpu/shader/execution/tanh.bin": "fe7a619d",
+  "webgpu/shader/execution/transpose.bin": "469edd7e",
+  "webgpu/shader/execution/trunc.bin": "8d3a05de",
+  "webgpu/shader/execution/unpack2x16float.bin": "e897c5ac",
+  "webgpu/shader/execution/unpack2x16snorm.bin": "450d5402",
+  "webgpu/shader/execution/unpack2x16unorm.bin": "306b3bf9",
+  "webgpu/shader/execution/unpack4x8snorm.bin": "fc1bd4c3",
+  "webgpu/shader/execution/unpack4x8unorm.bin": "763288cc",
+  "webgpu/shader/execution/unary/af_arithmetic.bin": "a39d4121",
+  "webgpu/shader/execution/unary/af_assignment.bin": "6ed540e8",
+  "webgpu/shader/execution/unary/bool_conversion.bin": "f37ea003",
+  "webgpu/shader/execution/unary/f16_arithmetic.bin": "9cadc739",
+  "webgpu/shader/execution/unary/f16_conversion.bin": "49c7b38e",
+  "webgpu/shader/execution/unary/f32_arithmetic.bin": "28da577a",
+  "webgpu/shader/execution/unary/f32_conversion.bin": "fcdfdd08",
+  "webgpu/shader/execution/unary/i32_arithmetic.bin": "25cf27d1",
+  "webgpu/shader/execution/unary/i32_conversion.bin": "276dcf69",
+  "webgpu/shader/execution/unary/u32_conversion.bin": "acac2172",
+  "webgpu/shader/execution/unary/ai_assignment.bin": "c876d431",
+  "webgpu/shader/execution/binary/ai_arithmetic.bin": "36a1d65b",
+  "webgpu/shader/execution/unary/ai_arithmetic.bin": "89e34dcf",
+  "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "3fd8797b",
+  "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "bf99158a",
+  "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "8bcc7c30"
 }
\ No newline at end of file
diff --git a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_addition.bin b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_addition.bin
index 634f518e0d12..ba9d123cbf41 100644
Binary files a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_addition.bin and b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_addition.bin differ
diff --git a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin
index 5f8636b77c7d..b710ae908402 100644
Binary files a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin and b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin differ
diff --git a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin
index a8d35f964d4e..c8a3b7205aaa 100644
Binary files a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin and b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin differ
diff --git a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_subtraction.bin b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_subtraction.bin
index 1754d77f1aad..8f88d196cebe 100644
Binary files a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_subtraction.bin and b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_subtraction.bin differ
diff --git a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin
index 0a61c01c1a46..e57f91ffd980 100644
Binary files a/src/resources/cache/webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin and b/src/resources/cache/webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin differ
diff --git a/src/unittests/floating_point.spec.ts b/src/unittests/floating_point.spec.ts
index 472207d185a1..0c5c32ecd2e3 100644
--- a/src/unittests/floating_point.spec.ts
+++ b/src/unittests/floating_point.spec.ts
@@ -1836,27 +1836,24 @@ interface AbsoluteErrorCase {
 const kSmallAbsoluteErrorValue = {
   f32: 2 ** -11, // Builtin cos and sin has a absolute error 2**-11 for f32
   f16: 2 ** -7, // Builtin cos and sin has a absolute error 2**-7 for f16
-  abstract: 2 ** -11, // Builtin cos and sin has a absolute error 2**-11 for AbstractFloat
 } as const;
 // A large absolute error value is a representable value x that much smaller than maximum
 // positive, but positive.max - x is still exactly representable.
 const kLargeAbsoluteErrorValue = {
   f32: 2 ** 110, // f32.positive.max - 2**110 = 3.4028104e+38 = 0x7f7fffbf in f32
   f16: 2 ** 10, // f16.positive.max - 2**10 = 64480 = 0x7bdf in f16
-  abstract: 2 ** 977, // f64.positive.man - 2**977 = 1.79769e+308 = 0x7fefffffffffffbf in f64
 } as const;
 // A subnormal absolute error value is a subnormal representable value x of kind, which ensures
 // that positive.subnormal.min +/- x is still exactly representable.
 const kSubnormalAbsoluteErrorValue = {
   f32: 2 ** -140, // f32 0x00000200
   f16: 2 ** -20, // f16 0x0010
-  abstract: 2 ** -1065, // f64 0x0000_0000_0000_0200
 } as const;
 
 g.test('absoluteErrorInterval')
   .params(u =>
     u
-      .combine('trait', ['f32', 'f16', 'abstract'] as const)
+      .combine('trait', ['f32', 'f16'] as const)
       .beginSubcases()
       .expandWithParams<AbsoluteErrorCase>(p => {
         const trait = FP[p.trait];
@@ -1865,26 +1862,6 @@ g.test('absoluteErrorInterval')
         const largeErr = kLargeAbsoluteErrorValue[p.trait];
         const subnormalErr = kSubnormalAbsoluteErrorValue[p.trait];
 
-        // Additional testing for non-f64 values, since JS number is f64 internally
-        // prettier-ignore
-        const additionalSubnormal64bit = p.trait !== 'abstract' ?
-          [
-            // 64-bit subnormals, expected to be treated as 0.0 or smallest subnormal of kind.
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: 0, expected: [0, constants.positive.subnormal.min] },
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: subnormalErr, expected: [-subnormalErr, constants.positive.subnormal.min + subnormalErr] },
-            // Note that f32 minimum subnormal is so smaller than 1.0, adding them together may result in the f64 results 1.0.
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: 1, expected: [-1, constants.positive.subnormal.min + 1] },
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: 0, expected: [0, constants.positive.subnormal.min] },
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: subnormalErr, expected: [-subnormalErr, constants.positive.subnormal.min + subnormalErr] },
-            { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: 1, expected: [-1, constants.positive.subnormal.min + 1] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: 0, expected: [constants.negative.subnormal.max, 0] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: subnormalErr, expected: [constants.negative.subnormal.max - subnormalErr, subnormalErr] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: 1, expected: [constants.negative.subnormal.max - 1, 1] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: 0, expected: [constants.negative.subnormal.max, 0] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: subnormalErr, expected: [constants.negative.subnormal.max - subnormalErr, subnormalErr] },
-            { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: 1, expected: [constants.negative.subnormal.max - 1, 1] },
-          ] as const : [] as const;
-
         // prettier-ignore
         return [
           // Edge Cases
@@ -1940,7 +1917,20 @@ g.test('absoluteErrorInterval')
           { value: -2, error: smallErr, expected: [-2 - smallErr, -2 + smallErr] },
           { value: -2, error: 1, expected: [-3, -1] },
 
-          ...additionalSubnormal64bit,
+          // 64-bit subnormals, expected to be treated as 0.0 or smallest subnormal of kind.
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: 0, expected: [0, constants.positive.subnormal.min] },
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: subnormalErr, expected: [-subnormalErr, constants.positive.subnormal.min + subnormalErr] },
+          // Note that f32 minimum subnormal is so smaller than 1.0, adding them together may result in the f64 results 1.0.
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0001n), error: 1, expected: [-1, constants.positive.subnormal.min + 1] },
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: 0, expected: [0, constants.positive.subnormal.min] },
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: subnormalErr, expected: [-subnormalErr, constants.positive.subnormal.min + subnormalErr] },
+          { value: reinterpretU64AsF64(0x0000_0000_0000_0002n), error: 1, expected: [-1, constants.positive.subnormal.min + 1] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: 0, expected: [constants.negative.subnormal.max, 0] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: subnormalErr, expected: [constants.negative.subnormal.max - subnormalErr, subnormalErr] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_ffffn), error: 1, expected: [constants.negative.subnormal.max - 1, 1] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: 0, expected: [constants.negative.subnormal.max, 0] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: subnormalErr, expected: [constants.negative.subnormal.max - subnormalErr, subnormalErr] },
+          { value: reinterpretU64AsF64(0x800f_ffff_ffff_fffen), error: 1, expected: [constants.negative.subnormal.max - 1, 1] },
         ];
       })
   )
@@ -2102,7 +2092,7 @@ const kULPErrorValue = {
 g.test('ulpInterval')
   .params(u =>
     u
-      .combine('trait', ['abstract', 'f32', 'f16'] as const)
+      .combine('trait', ['f32', 'f16'] as const)
       .beginSubcases()
       .expandWithParams<ULPCase>(p => {
         const trait = kFPTraitForULP[p.trait];
diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index 1fa9fe582c4b..a074bafa2c3c 100644
--- a/src/webgpu/listing_meta.json
+++ b/src/webgpu/listing_meta.json
@@ -1893,7 +1893,9 @@
   "webgpu:shader,validation,expression,call,builtin,exp2:values:*": { "subcaseMS": 0.410 },
   "webgpu:shader,validation,expression,call,builtin,exp:integer_argument:*": { "subcaseMS": 1.356 },
   "webgpu:shader,validation,expression,call,builtin,exp:values:*": { "subcaseMS": 0.311 },
+  "webgpu:shader,validation,expression,call,builtin,floor:bad_args:*": { "subcaseMS": 14.467 },
   "webgpu:shader,validation,expression,call,builtin,floor:integer_argument:*": { "subcaseMS": 48.400 },
+  "webgpu:shader,validation,expression,call,builtin,floor:must_use:*": { "subcaseMS": 0.170 },
   "webgpu:shader,validation,expression,call,builtin,floor:values:*": { "subcaseMS": 29.668 },
   "webgpu:shader,validation,expression,call,builtin,inverseSqrt:integer_argument:*": { "subcaseMS": 1.356 },
   "webgpu:shader,validation,expression,call,builtin,inverseSqrt:values:*": { "subcaseMS": 0.315 },
@@ -1940,6 +1942,10 @@
   "webgpu:shader,validation,expression,call,builtin,sqrt:values:*": { "subcaseMS": 0.302 },
   "webgpu:shader,validation,expression,call,builtin,tan:integer_argument:*": { "subcaseMS": 1.734 },
   "webgpu:shader,validation,expression,call,builtin,tan:values:*": { "subcaseMS": 0.350 },
+  "webgpu:shader,validation,expression,call,builtin,textureSample:array_index_argument:*": { "subcaseMS": 1.888 },
+  "webgpu:shader,validation,expression,call,builtin,textureSample:coords_argument:*": { "subcaseMS": 1.342 },
+  "webgpu:shader,validation,expression,call,builtin,textureSample:offset_argument,non_const:*": { "subcaseMS": 1.604 },
+  "webgpu:shader,validation,expression,call,builtin,textureSample:offset_argument:*": { "subcaseMS": 1.401 },
   "webgpu:shader,validation,expression,call,builtin,unpack4xI8:bad_args:*": { "subcaseMS": 121.263 },
   "webgpu:shader,validation,expression,call,builtin,unpack4xI8:must_use:*": { "subcaseMS": 35.200 },
   "webgpu:shader,validation,expression,call,builtin,unpack4xI8:supported:*": { "subcaseMS": 24.150 },
@@ -2177,6 +2183,9 @@
   "webgpu:shader,validation,types,struct:no_indirect_recursion_via_array_size:*": { "subcaseMS": 0.900 },
   "webgpu:shader,validation,types,struct:no_indirect_recursion_via_struct_attribute:*": { "subcaseMS": 1.467 },
   "webgpu:shader,validation,types,struct:no_indirect_recursion_via_struct_member_nested_in_alias:*": { "subcaseMS": 0.950 },
+  "webgpu:shader,validation,types,textures:storage_texture_types:*": { "subcaseMS": 167.510 },
+  "webgpu:shader,validation,types,textures:texel_formats,as_value:*": { "subcaseMS": 0.518 },
+  "webgpu:shader,validation,types,textures:texel_formats:*": { "subcaseMS": 1707.432 },
   "webgpu:shader,validation,types,vector:vector:*": { "subcaseMS": 1.295 },
   "webgpu:shader,validation,uniformity,uniformity:basics:*": { "subcaseMS": 1.467 },
   "webgpu:shader,validation,uniformity,uniformity:binary_expressions:*": { "subcaseMS": 1.758 },
diff --git a/src/webgpu/shader/execution/expression/binary/af_matrix_matrix_multiplication.cache.ts b/src/webgpu/shader/execution/expression/binary/af_matrix_matrix_multiplication.cache.ts
index c7540083d6ff..e539ee1db5da 100644
--- a/src/webgpu/shader/execution/expression/binary/af_matrix_matrix_multiplication.cache.ts
+++ b/src/webgpu/shader/execution/expression/binary/af_matrix_matrix_multiplication.cache.ts
@@ -11,7 +11,7 @@ const mat_mat_cases = ([2, 3, 4] as const)
         [`mat${k}x${rows}_mat${cols}x${k}`]: () => {
           return selectNCases(
             'binary/af_matrix_matrix_multiplication',
-            50,
+            10,
             FP.abstract.generateMatrixPairToMatrixCases(
               sparseMatrixF64Range(k, rows),
               sparseMatrixF64Range(cols, k),
diff --git a/src/webgpu/shader/execution/expression/case.ts b/src/webgpu/shader/execution/expression/case.ts
index 063b58da4267..d837b1c32fb4 100644
--- a/src/webgpu/shader/execution/expression/case.ts
+++ b/src/webgpu/shader/execution/expression/case.ts
@@ -69,7 +69,9 @@ export function selectNCases(dis: string, n: number, cases: Case[]): Case[] {
     return cases;
   }
   const dis_crc32 = crc32(dis);
-  return cases.filter(c => n * (0xffff_ffff / count) > (crc32(c.input.toString()) ^ dis_crc32));
+  return cases.filter(
+    c => Math.trunc((n / count) * 0xffff_ffff) > (crc32(c.input.toString()) ^ dis_crc32) >>> 0
+  );
 }
 
 /**
diff --git a/src/webgpu/shader/execution/expression/expression.ts b/src/webgpu/shader/execution/expression/expression.ts
index 50885c0d3594..e86d068c6daa 100644
--- a/src/webgpu/shader/execution/expression/expression.ts
+++ b/src/webgpu/shader/execution/expression/expression.ts
@@ -14,6 +14,7 @@ import {
   isAbstractType,
   scalarTypeOf,
   ArrayType,
+  elementTypeOf,
 } from '../../../util/conversion.js';
 import { align } from '../../../util/math.js';
 
@@ -157,7 +158,10 @@ function wgslMembers(members: Type[], source: InputSource, memberName: (i: numbe
   });
   const padding = layout.stride - layout.size;
   if (padding > 0) {
-    lines.push(`  @size(${padding}) padding : i32,`);
+    // Pad with a 'f16' if the padding requires an odd multiple of 2 bytes.
+    // This is required as 'i32' has an alignment and size of 4 bytes.
+    const ty = (padding & 2) !== 0 ? 'f16' : 'i32';
+    lines.push(`  @size(${padding}) padding : ${ty},`);
   }
   return lines.join('\n');
 }
@@ -618,7 +622,7 @@ function basicExpressionShaderBody(
     // Constant eval
     //////////////////////////////////////////////////////////////////////////
     let body = '';
-    if (parameterTypes.some(isAbstractType)) {
+    if (parameterTypes.some(ty => isAbstractType(elementTypeOf(ty)))) {
       // Directly assign the expression to the output, to avoid an
       // intermediate store, which will concretize the value early
       body = cases
diff --git a/src/webgpu/shader/validation/expression/call/builtin/dot4I8Packed.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/dot4I8Packed.spec.ts
index 12a73ae0c602..56c5cf5403eb 100644
--- a/src/webgpu/shader/validation/expression/call/builtin/dot4I8Packed.spec.ts
+++ b/src/webgpu/shader/validation/expression/call/builtin/dot4I8Packed.spec.ts
@@ -19,6 +19,8 @@ const kBadArgs = {
   '1f32': '(1u,2f)',
   '1bool': '(1u,true)',
   '1vec2u': '(1u,vec2u())',
+  bool_bool: '(false,true)',
+  bool2_bool2: '(vec2<bool>(),vec2(false,true))',
 };
 
 export const g = makeTestGroup(ShaderValidationTest);
diff --git a/src/webgpu/shader/validation/expression/call/builtin/dot4U8Packed.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/dot4U8Packed.spec.ts
index 18fb97707da1..1d240af5a81b 100644
--- a/src/webgpu/shader/validation/expression/call/builtin/dot4U8Packed.spec.ts
+++ b/src/webgpu/shader/validation/expression/call/builtin/dot4U8Packed.spec.ts
@@ -19,6 +19,8 @@ const kBadArgs = {
   '1f32': '(1u,2f)',
   '1bool': '(1u,true)',
   '1vec2u': '(1u,vec2u())',
+  bool_bool: '(false,true)',
+  bool2_bool2: '(vec2<bool>(),vec2(false,true))',
 };
 
 export const g = makeTestGroup(ShaderValidationTest);
diff --git a/src/webgpu/shader/validation/expression/call/builtin/floor.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/floor.spec.ts
index b65fafc74466..e43c6fa3fe43 100644
--- a/src/webgpu/shader/validation/expression/call/builtin/floor.spec.ts
+++ b/src/webgpu/shader/validation/expression/call/builtin/floor.spec.ts
@@ -73,3 +73,32 @@ Validates that scalar and vector integer arguments are rejected by ${builtin}()
       'constant'
     );
   });
+
+const kGoodArgs = '(1.1)';
+const kBadArgs = {
+  no_parens: '',
+  // Bad number of args
+  '0args': '()',
+  '2args': '(1.0,2.0)',
+  // Bad value for arg 0
+  '0aint': '(1)',
+  '0i32': '(1i)',
+  '0u32': '(1u)',
+  '0bool': '(false)',
+  '0vec2u': '(vec2u())',
+  '0array': '(array(1.1,2.2))',
+  '0struct': '(modf(2.2))',
+};
+
+g.test('bad_args')
+  .desc(`Test compilation failure of ${builtin} with bad arguments`)
+  .params(u => u.combine('arg', keysOf(kBadArgs)))
+  .fn(t => {
+    t.expectCompileResult(false, `const c = ${builtin}${kBadArgs[t.params.arg]};`);
+  });
+
+g.test('must_use')
+  .desc(`Result of ${builtin} must be used`)
+  .fn(t => {
+    t.expectCompileResult(false, `fn f() { ${builtin}${kGoodArgs}; }`);
+  });
diff --git a/src/webgpu/shader/validation/expression/call/builtin/textureSample.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/textureSample.spec.ts
new file mode 100644
index 000000000000..0d63c059d719
--- /dev/null
+++ b/src/webgpu/shader/validation/expression/call/builtin/textureSample.spec.ts
@@ -0,0 +1,218 @@
+const builtin = 'textureSample';
+export const description = `
+Validation tests for the ${builtin}() builtin.
+
+* test textureSample coords parameter must be correct type
+* test textureSample array_index parameter must be correct type
+* test textureSample coords parameter must be correct type
+* test textureSample offset parameter must be correct type
+* test textureSample offset parameter must be a const-expression
+* test textureSample offset parameter must be between -8 and +7 inclusive
+`;
+
+import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
+import {
+  Type,
+  kAllScalarsAndVectors,
+  isConvertible,
+  ScalarType,
+  VectorType,
+} from '../../../../../util/conversion.js';
+import { ShaderValidationTest } from '../../../shader_validation_test.js';
+
+type TextureSampleArguments = {
+  coordsArgType: ScalarType | VectorType;
+  hasArrayIndexArg?: boolean;
+  offsetArgType?: VectorType;
+};
+
+const kValidTextureSampleParameterTypes: { [n: string]: TextureSampleArguments } = {
+  'texture_1d<f32>': { coordsArgType: Type.f32 },
+  'texture_2d<f32>': { coordsArgType: Type.vec2f, offsetArgType: Type.vec2i },
+  'texture_2d_array<f32>': {
+    coordsArgType: Type.vec2f,
+    hasArrayIndexArg: true,
+    offsetArgType: Type.vec2i,
+  },
+  'texture_3d<f32>': { coordsArgType: Type.vec3f, offsetArgType: Type.vec3i },
+  'texture_cube<f32>': { coordsArgType: Type.vec3f },
+  'texture_cube_array<f32>': { coordsArgType: Type.vec3f, hasArrayIndexArg: true },
+  texture_depth_2d: { coordsArgType: Type.vec2f, offsetArgType: Type.vec2i },
+  texture_depth_2d_array: { coordsArgType: Type.vec2f, hasArrayIndexArg: true },
+  texture_depth_cube: { coordsArgType: Type.vec3f },
+  texture_depth_cube_array: { coordsArgType: Type.vec3f, hasArrayIndexArg: true },
+} as const;
+
+const kTextureTypes = keysOf(kValidTextureSampleParameterTypes);
+const kValuesTypes = objectsToRecord(kAllScalarsAndVectors);
+
+export const g = makeTestGroup(ShaderValidationTest);
+
+g.test('coords_argument')
+  .desc(
+    `
+Validates that only incorrect coords arguments are rejected by ${builtin}
+`
+  )
+  .params(u =>
+    u
+      .combine('textureType', keysOf(kValidTextureSampleParameterTypes))
+      .combine('coordType', keysOf(kValuesTypes))
+      .beginSubcases()
+      .combine('value', [-1, 0, 1] as const)
+      .expand('offset', ({ textureType }) => {
+        const offset = kValidTextureSampleParameterTypes[textureType].offsetArgType;
+        return offset ? [false, true] : [false];
+      })
+  )
+  .fn(t => {
+    const { textureType, coordType, offset, value } = t.params;
+    const coordArgType = kValuesTypes[coordType];
+    const {
+      offsetArgType,
+      coordsArgType: coordsRequiredType,
+      hasArrayIndexArg,
+    } = kValidTextureSampleParameterTypes[textureType];
+
+    const coordWGSL = coordArgType.create(value).wgsl();
+    const arrayWGSL = hasArrayIndexArg ? ', 0' : '';
+    const offsetWGSL = offset ? `, ${offsetArgType?.create(0).wgsl()}` : '';
+
+    const code = `
+@group(0) @binding(0) var s: sampler;
+@group(0) @binding(1) var t: ${textureType};
+@fragment fn fs() -> @location(0) vec4f {
+  let v = textureSample(t, s, ${coordWGSL}${arrayWGSL}${offsetWGSL});
+  return vec4f(0);
+}
+`;
+    const expectSuccess = isConvertible(coordArgType, coordsRequiredType);
+    t.expectCompileResult(expectSuccess, code);
+  });
+
+g.test('array_index_argument')
+  .desc(
+    `
+Validates that only incorrect array_index arguments are rejected by ${builtin}
+`
+  )
+  .params(u =>
+    u
+      .combine('textureType', kTextureTypes)
+      // filter out types with no array_index
+      .filter(
+        ({ textureType }) => !!kValidTextureSampleParameterTypes[textureType].hasArrayIndexArg
+      )
+      .combine('arrayIndexType', keysOf(kValuesTypes))
+      .beginSubcases()
+      .combine('value', [-9, -8, 0, 7, 8])
+  )
+  .fn(t => {
+    const { textureType, arrayIndexType, value } = t.params;
+    const arrayIndexArgType = kValuesTypes[arrayIndexType];
+    const args = [arrayIndexArgType.create(value)];
+    const { coordsArgType, offsetArgType } = kValidTextureSampleParameterTypes[textureType];
+
+    const coordWGSL = coordsArgType.create(0).wgsl();
+    const arrayWGSL = args.map(arg => arg.wgsl()).join(', ');
+    const offsetWGSL = offsetArgType ? `, ${offsetArgType.create(0).wgsl()}` : '';
+
+    const code = `
+@group(0) @binding(0) var s: sampler;
+@group(0) @binding(1) var t: ${textureType};
+@fragment fn fs() -> @location(0) vec4f {
+  let v = textureSample(t, s, ${coordWGSL}, ${arrayWGSL}${offsetWGSL});
+  return vec4f(0);
+}
+`;
+    const expectSuccess =
+      isConvertible(arrayIndexArgType, Type.i32) || isConvertible(arrayIndexArgType, Type.u32);
+    t.expectCompileResult(expectSuccess, code);
+  });
+
+g.test('offset_argument')
+  .desc(
+    `
+Validates that only incorrect offset arguments are rejected by ${builtin}
+`
+  )
+  .params(u =>
+    u
+      .combine('textureType', kTextureTypes)
+      // filter out types with no offset
+      .filter(
+        ({ textureType }) =>
+          kValidTextureSampleParameterTypes[textureType].offsetArgType !== undefined
+      )
+      .combine('offsetType', keysOf(kValuesTypes))
+      .beginSubcases()
+      .combine('value', [-9, -8, 0, 7, 8])
+  )
+  .fn(t => {
+    const { textureType, offsetType, value } = t.params;
+    const offsetArgType = kValuesTypes[offsetType];
+    const args = [offsetArgType.create(value)];
+    const {
+      coordsArgType,
+      hasArrayIndexArg,
+      offsetArgType: offsetRequiredType,
+    } = kValidTextureSampleParameterTypes[textureType];
+
+    const coordWGSL = coordsArgType.create(0).wgsl();
+    const arrayWGSL = hasArrayIndexArg ? ', 0' : '';
+    const offsetWGSL = args.map(arg => arg.wgsl()).join(', ');
+
+    const code = `
+@group(0) @binding(0) var s: sampler;
+@group(0) @binding(1) var t: ${textureType};
+@fragment fn fs() -> @location(0) vec4f {
+  let v = textureSample(t, s, ${coordWGSL}${arrayWGSL}, ${offsetWGSL});
+  return vec4f(0);
+}
+`;
+    const expectSuccess =
+      isConvertible(offsetArgType, offsetRequiredType!) && value >= -8 && value <= 7;
+    t.expectCompileResult(expectSuccess, code);
+  });
+
+g.test('offset_argument,non_const')
+  .desc(
+    `
+Validates that only non-const offset arguments are rejected by ${builtin}
+`
+  )
+  .params(u =>
+    u
+      .combine('textureType', kTextureTypes)
+      .combine('varType', ['c', 'u', 'l'])
+      // filter out types with no offset
+      .filter(
+        ({ textureType }) =>
+          kValidTextureSampleParameterTypes[textureType].offsetArgType !== undefined
+      )
+  )
+  .fn(t => {
+    const { textureType, varType } = t.params;
+    const { coordsArgType, hasArrayIndexArg, offsetArgType } =
+      kValidTextureSampleParameterTypes[textureType];
+
+    const coordWGSL = coordsArgType.create(0).wgsl();
+    const arrayWGSL = hasArrayIndexArg ? ', 0' : '';
+    const offsetWGSL = `${offsetArgType}(${varType})`;
+    const castWGSL = offsetArgType!.elementType.toString();
+
+    const code = `
+@group(0) @binding(0) var s: sampler;
+@group(0) @binding(1) var t: ${textureType};
+@group(0) @binding(2) var<uniform> u: ${offsetArgType};
+@fragment fn fs(@builtin(position) p: vec4f) -> @location(0) vec4f {
+  const c = 1;
+  let l = ${offsetArgType}(${castWGSL}(p.x));
+  let v = textureSample(t, s, ${coordWGSL}${arrayWGSL}, ${offsetWGSL});
+  return vec4f(0);
+}
+`;
+    const expectSuccess = varType === 'c';
+    t.expectCompileResult(expectSuccess, code);
+  });
diff --git a/src/webgpu/shader/validation/types/textures.spec.ts b/src/webgpu/shader/validation/types/textures.spec.ts
new file mode 100644
index 000000000000..904fa4721b8d
--- /dev/null
+++ b/src/webgpu/shader/validation/types/textures.spec.ts
@@ -0,0 +1,92 @@
+export const description = `
+Validation tests for various texture types in shaders.
+
+TODO:
+- Sampled Texture Types
+- Multisampled Texture Types
+- External Sampled Texture Types
+- Depth Texture Types
+- Sampler Types
+`;
+
+import { makeTestGroup } from '../../../../common/framework/test_group.js';
+import {
+  isTextureFormatUsableAsStorageFormat,
+  kAllTextureFormats,
+  kColorTextureFormats,
+  kTextureFormatInfo,
+} from '../../../format_info.js';
+import { getPlainTypeInfo } from '../../../util/shader.js';
+import { ShaderValidationTest } from '../shader_validation_test.js';
+
+export const g = makeTestGroup(ShaderValidationTest);
+
+g.test('texel_formats')
+  .desc(
+    'Test channels and channel format of various texel formats when using as the storage texture format'
+  )
+  .params(u =>
+    u
+      .combine('format', kColorTextureFormats)
+      .filter(p => kTextureFormatInfo[p.format].color.storage)
+      .beginSubcases()
+      .combine('shaderScalarType', ['f32', 'u32', 'i32', 'bool', 'f16'] as const)
+  )
+  .beforeAllSubcases(t => {
+    if (t.params.shaderScalarType === 'f16') {
+      t.selectDeviceOrSkipTestCase({ requiredFeatures: ['shader-f16'] });
+    }
+
+    if (!isTextureFormatUsableAsStorageFormat(t.params.format, t.isCompatibility)) {
+      t.skip('storage usage is unsupported');
+    }
+  })
+  .fn(t => {
+    const { format, shaderScalarType } = t.params;
+    const info = kTextureFormatInfo[format];
+    const validShaderScalarType = getPlainTypeInfo(info.color.type);
+    const shaderValueType = `vec4<${shaderScalarType}>`;
+    const wgsl = `
+    @group(0) @binding(0) var tex: texture_storage_2d<${format}, read>;
+    @compute @workgroup_size(1) fn main() {
+        let v : ${shaderValueType} = textureLoad(tex, vec2u(0));
+        _ = v;
+    }
+`;
+    t.expectCompileResult(validShaderScalarType === shaderScalarType, wgsl);
+  });
+
+g.test('texel_formats,as_value')
+  .desc('Test that texel format cannot be used as value')
+  .fn(t => {
+    const wgsl = `
+    @compute @workgroup_size(1) fn main() {
+        var i = rgba8unorm;
+    }
+`;
+    t.expectCompileResult(false, wgsl);
+  });
+
+const kAccessModes = ['read', 'write', 'read_write'];
+
+g.test('storage_texture_types')
+  .desc(
+    `Test that for texture_storage_xx<format, access>
+- format must be an enumerant for one of the texel formats for storage textures
+- access must be an enumerant for one of the access modes
+
+Besides, the shader compilation should always pass regardless of whether the format supports the usage indicated by the access or not.
+`
+  )
+  .params(u =>
+    u.combine('access', [...kAccessModes, 'storage'] as const).combine('format', kAllTextureFormats)
+  )
+  .fn(t => {
+    const { format, access } = t.params;
+    const info = kTextureFormatInfo[format];
+    const isFormatValid =
+      info.color?.storage || info.depth?.storage || info.stencil?.storage || false;
+    const isAccessValid = kAccessModes.includes(access);
+    const wgsl = `@group(0) @binding(0) var tex: texture_storage_2d<${format}, ${access}>;`;
+    t.expectCompileResult(isFormatValid && isAccessValid, wgsl);
+  });
diff --git a/src/webgpu/util/conversion.ts b/src/webgpu/util/conversion.ts
index 15557b3d4c57..4925899811b5 100644
--- a/src/webgpu/util/conversion.ts
+++ b/src/webgpu/util/conversion.ts
@@ -2234,6 +2234,56 @@ export function isFloatType(ty: Type): boolean {
   return false;
 }
 
+/** @returns true if an argument of type 'src' can be used for a parameter of type 'dst' */
+export function isConvertible(src: Type, dst: Type) {
+  if (src === dst) {
+    return true;
+  }
+
+  const widthOf = (ty: Type) => {
+    return ty instanceof VectorType ? ty.width : 1;
+  };
+
+  if (widthOf(src) !== widthOf(dst)) {
+    return false;
+  }
+
+  const elSrc = scalarTypeOf(src);
+  const elDst = scalarTypeOf(dst);
+
+  switch (elSrc.kind) {
+    case 'abstract-float':
+      switch (elDst.kind) {
+        case 'abstract-float':
+        case 'f16':
+        case 'f32':
+        case 'f64':
+          return true;
+        default:
+          return false;
+      }
+    case 'abstract-int':
+      switch (elDst.kind) {
+        case 'abstract-int':
+        case 'abstract-float':
+        case 'f16':
+        case 'f32':
+        case 'f64':
+        case 'u16':
+        case 'u32':
+        case 'u8':
+        case 'i16':
+        case 'i32':
+        case 'i8':
+          return true;
+        default:
+          return false;
+      }
+    default:
+      return false;
+  }
+}
+
 /// All floating-point scalar types
 const kFloatScalars = [Type.abstractFloat, Type.f32, Type.f16] as const;
 
diff --git a/src/webgpu/util/floating_point.ts b/src/webgpu/util/floating_point.ts
index 36900e6da994..7887d82aaf70 100644
--- a/src/webgpu/util/floating_point.ts
+++ b/src/webgpu/util/floating_point.ts
@@ -1047,14 +1047,14 @@ export abstract class FPTraits {
     unreachable(`'refract' is not yet implemented for '${this.kind}'`);
   }
 
-  /** Version of absoluteErrorInterval that always returns the unboundedInterval */
-  protected unboundedAbsoluteErrorInterval(_n: number, _error_range: number): FPInterval {
-    return this.constants().unboundedInterval;
+  /** Stub for absolute errors */
+  protected unimplementedAbsoluteErrorInterval(_n: number, _error_range: number): FPInterval {
+    unreachable(`Absolute Error is not implement for '${this.kind}'`);
   }
 
-  /** Version of ulpInterval that always returns the unboundedInterval */
-  protected unboundedUlpInterval(_n: number, _numULP: number): FPInterval {
-    return this.constants().unboundedInterval;
+  /** Stub for ULP errors */
+  protected unimplementedUlpInterval(_n: number, _numULP: number): FPInterval {
+    unreachable(`ULP Error is not implement for '${this.kind}'`);
   }
 
   // Utilities - Defined by subclass
@@ -5091,12 +5091,10 @@ class FPAbstractTraits extends FPTraits {
   public readonly sparseMatrixRange = sparseMatrixF64Range;
 
   // Framework - Fundamental Error Intervals - Overrides
-  public readonly absoluteErrorInterval = this.absoluteErrorIntervalImpl.bind(this);
+  public readonly absoluteErrorInterval = this.unimplementedAbsoluteErrorInterval.bind(this); // Should use FP.f32 instead
   public readonly correctlyRoundedInterval = this.correctlyRoundedIntervalImpl.bind(this);
   public readonly correctlyRoundedMatrix = this.correctlyRoundedMatrixImpl.bind(this);
-  public readonly ulpInterval = (n: number, numULP: number): FPInterval => {
-    return this.toInterval(kF32Traits.ulpInterval(n, numULP));
-  };
+  public readonly ulpInterval = this.unimplementedUlpInterval.bind(this); // Should use FP.f32 instead
 
   // Framework - API - Overrides
   public readonly absInterval = this.absIntervalImpl.bind(this);
diff --git a/src/webgpu/web_platform/worker/worker.spec.ts b/src/webgpu/web_platform/worker/worker.spec.ts
index deed5d75e232..14fe135633c9 100644
--- a/src/webgpu/web_platform/worker/worker.spec.ts
+++ b/src/webgpu/web_platform/worker/worker.spec.ts
@@ -1,9 +1,10 @@
 export const description = `
 Tests WebGPU is available in a dedicated worker and a shared 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.
+Note: Any CTS test can be run in a worker by passing ?worker=dedicated, ?worker=shared,
+?worker=service as a query parameter. The tests in this file are specifically to check
+that WebGPU is available in each worker type. When run in combination with a ?worker flag,
+they will test workers created from other workers (where APIs exist to do so).
 
 TODO[2]: Figure out how to make these tests run in service workers (not actually
 important unless service workers gain the ability to launch other workers).