diff --git a/src/common/internal/file_loader.ts b/src/common/internal/file_loader.ts
index b5e1b1a4460b..aae4b8799549 100644
--- a/src/common/internal/file_loader.ts
+++ b/src/common/internal/file_loader.ts
@@ -73,8 +73,9 @@ export abstract class TestFileLoader extends EventTarget {
     query: TestQuery,
     {
       subqueriesToExpand = [],
+      fullyExpandSubtrees = [],
       maxChunkTime = Infinity,
-    }: { subqueriesToExpand?: string[]; maxChunkTime?: number } = {}
+    }: { subqueriesToExpand?: string[]; fullyExpandSubtrees?: string[]; maxChunkTime?: number } = {}
   ): Promise<TestTree> {
     const tree = await loadTreeForQuery(this, query, {
       subqueriesToExpand: subqueriesToExpand.map(s => {
@@ -82,6 +83,7 @@ export abstract class TestFileLoader extends EventTarget {
         assert(q.level >= 2, () => `subqueriesToExpand entries should not be multi-file:\n  ${q}`);
         return q;
       }),
+      fullyExpandSubtrees: fullyExpandSubtrees.map(s => parseQuery(s)),
       maxChunkTime,
     });
     this.dispatchEvent(new MessageEvent<void>('finish'));
diff --git a/src/common/internal/tree.ts b/src/common/internal/tree.ts
index 594837059ca7..f2fad590373c 100644
--- a/src/common/internal/tree.ts
+++ b/src/common/internal/tree.ts
@@ -286,8 +286,9 @@ export async function loadTreeForQuery(
   queryToLoad: TestQuery,
   {
     subqueriesToExpand,
+    fullyExpandSubtrees = [],
     maxChunkTime = Infinity,
-  }: { subqueriesToExpand: TestQuery[]; maxChunkTime?: number }
+  }: { subqueriesToExpand: TestQuery[]; fullyExpandSubtrees?: TestQuery[]; maxChunkTime?: number }
 ): Promise<TestTree> {
   const suite = queryToLoad.suite;
   const specs = await loader.listing(suite);
@@ -303,6 +304,10 @@ export async function loadTreeForQuery(
       // If toExpand == subquery, no expansion is needed (but it's still "seen").
       if (ordering === Ordering.Equal) seenSubqueriesToExpand[i] = true;
       return ordering !== Ordering.StrictSubset;
+    }) &&
+    fullyExpandSubtrees.every(toExpand => {
+      const ordering = compareQueries(toExpand, subquery);
+      return ordering === Ordering.Unordered;
     });
 
   // L0 = suite-level, e.g. suite:*
diff --git a/src/common/runtime/helper/options.ts b/src/common/runtime/helper/options.ts
index 38974b803fac..a8b93974c74d 100644
--- a/src/common/runtime/helper/options.ts
+++ b/src/common/runtime/helper/options.ts
@@ -103,6 +103,30 @@ function getOptionsInfoFromSearchString<Type extends CTSOptions>(
   return optionValues as unknown as Type;
 }
 
+/**
+ * converts foo/bar/src/webgpu/this/that/file.spec.ts to webgpu:this,that,file,*
+ */
+function convertPathToQuery(path: string) {
+  // removes .spec.ts and splits by directory separators.
+  const parts = path.substring(0, path.length - 8).split(/\/|\\/g);
+  // Gets parts only after the last `src`. Example: returns ['webgpu', 'foo', 'bar', 'test']
+  // for ['Users', 'me', 'src', 'cts', 'src', 'webgpu', 'foo', 'bar', 'test']
+  const partsAfterSrc = parts.slice(parts.lastIndexOf('src') + 1);
+  const suite = partsAfterSrc.shift();
+  return `${suite}:${partsAfterSrc.join(',')},*`;
+}
+
+/**
+ * If a query looks like a path (ends in .spec.ts and has directory separators)
+ * then convert try to convert it to a query.
+ */
+function convertPathLikeToQuery(queryOrPath: string) {
+  return queryOrPath.endsWith('.spec.ts') &&
+    (queryOrPath.includes('/') || queryOrPath.includes('\\'))
+    ? convertPathToQuery(queryOrPath)
+    : queryOrPath;
+}
+
 /**
  * Given a test query string in the form of `suite:foo,bar,moo&opt1=val1&opt2=val2
  * returns the query and the options.
@@ -115,7 +139,7 @@ export function parseSearchParamLikeWithOptions<Type extends CTSOptions>(
   options: Type;
 } {
   const searchString = query.includes('q=') || query.startsWith('?') ? query : `q=${query}`;
-  const queries = new URLSearchParams(searchString).getAll('q');
+  const queries = new URLSearchParams(searchString).getAll('q').map(convertPathLikeToQuery);
   const options = getOptionsInfoFromSearchString(optionsInfos, searchString);
   return { queries, options };
 }
diff --git a/src/common/tools/gen_wpt_cts_html.ts b/src/common/tools/gen_wpt_cts_html.ts
index e8161304e90f..90c7cd4ef4c5 100644
--- a/src/common/tools/gen_wpt_cts_html.ts
+++ b/src/common/tools/gen_wpt_cts_html.ts
@@ -71,6 +71,13 @@ interface ConfigJSON {
     /** The prefix to trim from every line of the expectations_file. */
     prefix: string;
   };
+  /** Expend all subtrees for provided queries */
+  fullyExpandSubtrees?: {
+    file: string;
+    prefix: string;
+  };
+  /*No long path assert */
+  noLongPathAssert?: boolean;
 }
 
 interface Config {
@@ -79,10 +86,15 @@ interface Config {
   template: string;
   maxChunkTimeMS: number;
   argumentsPrefixes: string[];
+  noLongPathAssert: boolean;
   expectations?: {
     file: string;
     prefix: string;
   };
+  fullyExpandSubtrees?: {
+    file: string;
+    prefix: string;
+  };
 }
 
 let config: Config;
@@ -101,6 +113,7 @@ let config: Config;
         template: path.resolve(jsonFileDir, configJSON.template),
         maxChunkTimeMS: configJSON.maxChunkTimeMS ?? Infinity,
         argumentsPrefixes: configJSON.argumentsPrefixes ?? ['?q='],
+        noLongPathAssert: configJSON.noLongPathAssert ?? false,
       };
       if (configJSON.expectations) {
         config.expectations = {
@@ -108,6 +121,12 @@ let config: Config;
           prefix: configJSON.expectations.prefix,
         };
       }
+      if (configJSON.fullyExpandSubtrees) {
+        config.fullyExpandSubtrees = {
+          file: path.resolve(jsonFileDir, configJSON.fullyExpandSubtrees.file),
+          prefix: configJSON.fullyExpandSubtrees.prefix,
+        };
+      }
       break;
     }
     case 4:
@@ -130,6 +149,7 @@ let config: Config;
         suite,
         maxChunkTimeMS: Infinity,
         argumentsPrefixes: ['?q='],
+        noLongPathAssert: false,
       };
       if (process.argv.length >= 7) {
         config.argumentsPrefixes = (await fs.readFile(argsPrefixesFile, 'utf8'))
@@ -153,29 +173,16 @@ let config: Config;
   config.argumentsPrefixes.sort((a, b) => b.length - a.length);
 
   // Load expectations (if any)
-  let expectationLines = new Set<string>();
-  if (config.expectations) {
-    expectationLines = new Set(
-      (await fs.readFile(config.expectations.file, 'utf8')).split(/\r?\n/).filter(l => l.length)
-    );
-  }
+  const expectations: Map<string, string[]> = await loadQueryFile(
+    config.argumentsPrefixes,
+    config.expectations
+  );
 
-  const expectations: Map<string, string[]> = new Map();
-  for (const prefix of config.argumentsPrefixes) {
-    expectations.set(prefix, []);
-  }
-
-  expLoop: for (const exp of expectationLines) {
-    // Take each expectation for the longest prefix it matches.
-    for (const argsPrefix of config.argumentsPrefixes) {
-      const prefix = config.expectations!.prefix + argsPrefix;
-      if (exp.startsWith(prefix)) {
-        expectations.get(argsPrefix)!.push(exp.substring(prefix.length));
-        continue expLoop;
-      }
-    }
-    console.log('note: ignored expectation: ' + exp);
-  }
+  // Load fullyExpandSubtrees queries (if any)
+  const fullyExpand: Map<string, string[]> = await loadQueryFile(
+    config.argumentsPrefixes,
+    config.fullyExpandSubtrees
+  );
 
   const loader = new DefaultTestFileLoader();
   const lines = [];
@@ -183,6 +190,7 @@ let config: Config;
     const rootQuery = new TestQueryMultiFile(config.suite, []);
     const tree = await loader.loadTree(rootQuery, {
       subqueriesToExpand: expectations.get(prefix),
+      fullyExpandSubtrees: fullyExpand.get(prefix),
       maxChunkTime: config.maxChunkTimeMS,
     });
 
@@ -199,19 +207,21 @@ let config: Config;
       alwaysExpandThroughLevel,
     })) {
       assert(query instanceof TestQueryMultiCase);
-      const queryString = query.toString();
-      // Check for a safe-ish path length limit. Filename must be <= 255, and on Windows the whole
-      // path must be <= 259. Leave room for e.g.:
-      // 'c:\b\s\w\xxxxxxxx\layout-test-results\external\wpt\webgpu\cts_worker=0_q=...-actual.txt'
-      assert(
-        queryString.length < 185,
-        `Generated test variant would produce too-long -actual.txt filename. Possible solutions:
+      if (!config.noLongPathAssert) {
+        const queryString = query.toString();
+        // Check for a safe-ish path length limit. Filename must be <= 255, and on Windows the whole
+        // path must be <= 259. Leave room for e.g.:
+        // 'c:\b\s\w\xxxxxxxx\layout-test-results\external\wpt\webgpu\cts_worker=0_q=...-actual.txt'
+        assert(
+          queryString.length < 185,
+          `Generated test variant would produce too-long -actual.txt filename. Possible solutions:
 - Reduce the length of the parts of the test query
 - Reduce the parameterization of the test
 - Make the test function faster and regenerate the listing_meta entry
 - Reduce the specificity of test expectations (if you're using them)
 ${queryString}`
-      );
+        );
+      }
 
       lines.push({
         urlQueryString: prefix + query.toString(), // "?worker=0&q=..."
@@ -232,6 +242,39 @@ ${queryString}`
   process.exit(1);
 });
 
+async function loadQueryFile(
+  argumentsPrefixes: string[],
+  queryFile?: {
+    file: string;
+    prefix: string;
+  }
+): Promise<Map<string, string[]>> {
+  let lines = new Set<string>();
+  if (queryFile) {
+    lines = new Set(
+      (await fs.readFile(queryFile.file, 'utf8')).split(/\r?\n/).filter(l => l.length)
+    );
+  }
+
+  const result: Map<string, string[]> = new Map();
+  for (const prefix of argumentsPrefixes) {
+    result.set(prefix, []);
+  }
+
+  expLoop: for (const exp of lines) {
+    // Take each expectation for the longest prefix it matches.
+    for (const argsPrefix of argumentsPrefixes) {
+      const prefix = queryFile!.prefix + argsPrefix;
+      if (exp.startsWith(prefix)) {
+        result.get(argsPrefix)!.push(exp.substring(prefix.length));
+        continue expLoop;
+      }
+    }
+    console.log('note: ignored expectation: ' + exp);
+  }
+  return result;
+}
+
 async function generateFile(
   lines: Array<{ urlQueryString?: string; comment?: string } | undefined>
 ): Promise<void> {
diff --git a/src/resources/README.md b/src/resources/README.md
index daa51765baa0..375e5db55d00 100644
--- a/src/resources/README.md
+++ b/src/resources/README.md
@@ -10,7 +10,7 @@ ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libvpx -pix_fmt yuv420p -frames 50
 // Generate four-colors-theora-bt601.ogv, mimeType: 'video/ogg; codecs=theora'
 ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libtheora -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv four-colors-theora-bt601.ogv
 
-// Generate four-colors-h264-bt601.mp4, mimeType: 'video/mp4; codecs=h264'
+// Generate four-colors-h264-bt601.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
 ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv four-colors-h264-bt601.mp4
 
 // Generate four-colors-vp9-bt601.webm, mimeType: 'video/webm; codecs=vp9'
@@ -28,17 +28,17 @@ Use ffmepg to rotate video content x degrees in cw direction (by using `transpos
 
 H264 rotated video files are generated by ffmpeg cmds below:
 ```
-// Generate four-colors-h264-bt601-rotate-90.mp4, mimeType: 'video/mp4; codecs=h264'
+// Generate four-colors-h264-bt601-rotate-90.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
 ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv -vf transpose=2 temp.mp4
 ffmpeg -display_rotation 270 -i temp.mp4 -c copy four-colors-h264-bt601-rotate-90.mp4
 rm temp.mp4
 
-// Generate four-colors-h264-bt601-rotate-180.mp4, mimeType: 'video/mp4; codecs=h264'
+// Generate four-colors-h264-bt601-rotate-180.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
 ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv -vf transpose=2,transpose=2 temp.mp4
 ffmpeg -display_rotation 180 -i temp.mp4 -c copy four-colors-h264-bt601-rotate-180.mp4
 rm temp.mp4
 
-// Generate four-colors-h264-bt601-rotate-270.mp4, mimeType: 'video/mp4; codecs=h264'
+// Generate four-colors-h264-bt601-rotate-270.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
 ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv -vf transpose=1 temp.mp4
 ffmpeg -display_rotation 90 -i temp.mp4 -c copy four-colors-h264-bt601-rotate-270.mp4
 rm temp.mp4
@@ -62,4 +62,35 @@ ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libvpx-vp9 -pix_fmt yuv420p -frames
 ffmpeg -display_rotation 90 -i temp.mp4 -c copy four-colors-vp9-bt601-rotate-270.mp4
 rm temp.mp4
 
-```
\ No newline at end of file
+```
+
+Generate video files to test flip behaviour.
+Use ffmpeg to flip video content. Using `display_hflip` to do horizontal flip and `display_vflip` to do vertical flip.
+
+H264 flip video files are generated by ffmpeg cmds below:
+```
+// Generate four-colors-h264-bt601-hflip.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
+ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv temp.mp4
+ffmpeg -display_hflip -i temp.mp4 -c copy four-colors-h264-bt601-hflip.mp4
+rm temp.mp4
+
+// Generate four-colors-h264-bt601-vflip.mp4, mimeType: 'video/mp4; codecs=avc1.4d400c'
+ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libx264 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv temp.mp4
+ffmpeg -display_vflip -i temp.mp4 -c copy four-colors-h264-bt601-vflip.mp4
+rm temp.mp4
+
+```
+
+Vp9 flip video files are generated by ffmpeg cmds below:
+```
+// Generate four-colors-vp9-bt601-hflip.mp4, mimeType: 'video/mp4; codecs=vp09.00.10.08'
+ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libvpx-vp9 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv temp.mp4
+ffmpeg -display_hflip -i temp.mp4 -c copy four-colors-vp9-bt601-hflip.mp4
+rm temp.mp4
+
+// Generate four-colors-vp9-bt601-vflip.mp4, mimeType: 'video/mp4; codecs=vp09.00.10.08'
+ffmpeg.exe -loop 1 -i .\four-colors.png -c:v libvpx-vp9 -pix_fmt yuv420p -frames 50 -colorspace smpte170m -color_primaries smpte170m -color_trc smpte170m -color_range tv temp.mp4
+ffmpeg -display_vflip -i temp.mp4 -c copy four-colors-vp9-bt601-vflip.mp4
+rm temp.mp4
+
+```
diff --git a/src/resources/cache/hashes.json b/src/resources/cache/hashes.json
index d4c5bf049b7d..d3f1573af604 100644
--- a/src/resources/cache/hashes.json
+++ b/src/resources/cache/hashes.json
@@ -1,106 +1,106 @@
 {
-  "webgpu/shader/execution/binary/af_addition.bin": "a52dc67e",
-  "webgpu/shader/execution/binary/af_logical.bin": "27321b9c",
-  "webgpu/shader/execution/binary/af_division.bin": "27a6e445",
-  "webgpu/shader/execution/binary/af_matrix_addition.bin": "5baaa29c",
-  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "a6fa7d52",
-  "webgpu/shader/execution/binary/af_multiplication.bin": "6e00f4f5",
-  "webgpu/shader/execution/binary/af_remainder.bin": "ed7f7cb5",
-  "webgpu/shader/execution/binary/af_subtraction.bin": "943996c9",
-  "webgpu/shader/execution/binary/f16_addition.bin": "cfdc6eaf",
-  "webgpu/shader/execution/binary/f16_logical.bin": "431e624e",
-  "webgpu/shader/execution/binary/f16_division.bin": "3027fa33",
-  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "abfcd7e0",
-  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "d227e0c4",
-  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "5dbe55d6",
-  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "21055fd5",
-  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "618d746f",
-  "webgpu/shader/execution/binary/f16_multiplication.bin": "21c319a1",
-  "webgpu/shader/execution/binary/f16_remainder.bin": "3fbca362",
-  "webgpu/shader/execution/binary/f16_subtraction.bin": "d217dbdd",
-  "webgpu/shader/execution/binary/f32_addition.bin": "76aa3474",
-  "webgpu/shader/execution/binary/f32_logical.bin": "ff723c9d",
-  "webgpu/shader/execution/binary/f32_division.bin": "b9f8d082",
-  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "b2ab732e",
-  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "2aa2f18d",
-  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "5358d7b5",
-  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "d003d09b",
-  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "50c0c9d7",
-  "webgpu/shader/execution/binary/f32_multiplication.bin": "14f0ff95",
-  "webgpu/shader/execution/binary/f32_remainder.bin": "d8dda35",
-  "webgpu/shader/execution/binary/f32_subtraction.bin": "ec3197e1",
-  "webgpu/shader/execution/binary/i32_arithmetic.bin": "ae1c1d58",
-  "webgpu/shader/execution/binary/i32_comparison.bin": "46155b50",
-  "webgpu/shader/execution/binary/u32_arithmetic.bin": "bab5328e",
-  "webgpu/shader/execution/binary/u32_comparison.bin": "34d818e3",
-  "webgpu/shader/execution/abs.bin": "dfb8d72b",
-  "webgpu/shader/execution/acos.bin": "a7375dda",
-  "webgpu/shader/execution/acosh.bin": "1f174f89",
-  "webgpu/shader/execution/asin.bin": "c3f31628",
-  "webgpu/shader/execution/asinh.bin": "85f77cb4",
-  "webgpu/shader/execution/atan.bin": "fb903055",
-  "webgpu/shader/execution/atan2.bin": "9da751a4",
-  "webgpu/shader/execution/atanh.bin": "9d1b15d5",
-  "webgpu/shader/execution/bitcast.bin": "275beb25",
-  "webgpu/shader/execution/ceil.bin": "487bdc72",
-  "webgpu/shader/execution/clamp.bin": "b32ed8d1",
-  "webgpu/shader/execution/cos.bin": "3a441e50",
-  "webgpu/shader/execution/cosh.bin": "74616476",
-  "webgpu/shader/execution/cross.bin": "ba69174d",
-  "webgpu/shader/execution/degrees.bin": "d7635a8d",
-  "webgpu/shader/execution/determinant.bin": "7220ac6c",
-  "webgpu/shader/execution/distance.bin": "c2724ca0",
-  "webgpu/shader/execution/dot.bin": "2b9fb191",
-  "webgpu/shader/execution/exp.bin": "a208159e",
-  "webgpu/shader/execution/exp2.bin": "5bc4085d",
-  "webgpu/shader/execution/faceForward.bin": "5845d9d4",
-  "webgpu/shader/execution/floor.bin": "522bc246",
-  "webgpu/shader/execution/fma.bin": "762790f3",
-  "webgpu/shader/execution/fract.bin": "d02241aa",
-  "webgpu/shader/execution/frexp.bin": "8c29d0d0",
-  "webgpu/shader/execution/inverseSqrt.bin": "321dd6ab",
-  "webgpu/shader/execution/ldexp.bin": "b755835f",
-  "webgpu/shader/execution/length.bin": "8df127e6",
-  "webgpu/shader/execution/log.bin": "ec26f0f6",
-  "webgpu/shader/execution/log2.bin": "5032f9",
-  "webgpu/shader/execution/max.bin": "7e8c7c62",
-  "webgpu/shader/execution/min.bin": "785b98c2",
-  "webgpu/shader/execution/mix.bin": "148569aa",
-  "webgpu/shader/execution/modf.bin": "b071fff0",
-  "webgpu/shader/execution/normalize.bin": "776251fe",
-  "webgpu/shader/execution/pack2x16float.bin": "9c1bbb0",
-  "webgpu/shader/execution/pow.bin": "14adc127",
-  "webgpu/shader/execution/quantizeToF16.bin": "4d684508",
-  "webgpu/shader/execution/radians.bin": "fe6a106a",
-  "webgpu/shader/execution/reflect.bin": "991e7d90",
-  "webgpu/shader/execution/refract.bin": "9ce6c6e9",
-  "webgpu/shader/execution/round.bin": "2014bc85",
-  "webgpu/shader/execution/saturate.bin": "3c5b5d6a",
-  "webgpu/shader/execution/sign.bin": "5ca6e2c5",
-  "webgpu/shader/execution/sin.bin": "67f46a23",
-  "webgpu/shader/execution/sinh.bin": "85a3ebc6",
-  "webgpu/shader/execution/smoothstep.bin": "3650849c",
-  "webgpu/shader/execution/sqrt.bin": "199de8bb",
-  "webgpu/shader/execution/step.bin": "a1e0ecb1",
-  "webgpu/shader/execution/tan.bin": "48dfc5a4",
-  "webgpu/shader/execution/tanh.bin": "dba3d680",
-  "webgpu/shader/execution/transpose.bin": "8e5d0c34",
-  "webgpu/shader/execution/trunc.bin": "26bb567c",
-  "webgpu/shader/execution/unpack2x16float.bin": "dcbf83a8",
-  "webgpu/shader/execution/unpack2x16snorm.bin": "c85a9fdf",
-  "webgpu/shader/execution/unpack2x16unorm.bin": "7d8ce59f",
-  "webgpu/shader/execution/unpack4x8snorm.bin": "c394da0c",
-  "webgpu/shader/execution/unpack4x8unorm.bin": "e9849eb9",
-  "webgpu/shader/execution/unary/af_arithmetic.bin": "f5dd97a9",
-  "webgpu/shader/execution/unary/af_assignment.bin": "3dde81ac",
-  "webgpu/shader/execution/unary/bool_conversion.bin": "2b501a16",
-  "webgpu/shader/execution/unary/f16_arithmetic.bin": "e53cb569",
-  "webgpu/shader/execution/unary/f16_conversion.bin": "9bf49d51",
-  "webgpu/shader/execution/unary/f32_arithmetic.bin": "58d207c5",
-  "webgpu/shader/execution/unary/f32_conversion.bin": "56413b46",
-  "webgpu/shader/execution/unary/i32_arithmetic.bin": "8704047",
-  "webgpu/shader/execution/unary/i32_complement.bin": "7dec3502",
-  "webgpu/shader/execution/unary/i32_conversion.bin": "45acb16d",
-  "webgpu/shader/execution/unary/u32_complement.bin": "e000b062",
-  "webgpu/shader/execution/unary/u32_conversion.bin": "f2ffbc61"
+  "webgpu/shader/execution/binary/af_addition.bin": "cc376b37",
+  "webgpu/shader/execution/binary/af_logical.bin": "581be442",
+  "webgpu/shader/execution/binary/af_division.bin": "ca179ff2",
+  "webgpu/shader/execution/binary/af_matrix_addition.bin": "209e6965",
+  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "e1226295",
+  "webgpu/shader/execution/binary/af_multiplication.bin": "3c07556a",
+  "webgpu/shader/execution/binary/af_remainder.bin": "2a4fd923",
+  "webgpu/shader/execution/binary/af_subtraction.bin": "74fafb7a",
+  "webgpu/shader/execution/binary/f16_addition.bin": "54274ca7",
+  "webgpu/shader/execution/binary/f16_logical.bin": "b6e1c3f7",
+  "webgpu/shader/execution/binary/f16_division.bin": "b0798477",
+  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "9efe8261",
+  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "3b8fe08a",
+  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "495f7d20",
+  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "738b53d5",
+  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "66dea74b",
+  "webgpu/shader/execution/binary/f16_multiplication.bin": "fc8cf78",
+  "webgpu/shader/execution/binary/f16_remainder.bin": "da96b25a",
+  "webgpu/shader/execution/binary/f16_subtraction.bin": "58a68e4a",
+  "webgpu/shader/execution/binary/f32_addition.bin": "4b64583c",
+  "webgpu/shader/execution/binary/f32_logical.bin": "43089982",
+  "webgpu/shader/execution/binary/f32_division.bin": "3dbe6d9f",
+  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "dd8ff303",
+  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "d51044f4",
+  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "206f8b97",
+  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "216f822b",
+  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "60c1da25",
+  "webgpu/shader/execution/binary/f32_multiplication.bin": "55439f76",
+  "webgpu/shader/execution/binary/f32_remainder.bin": "962ef47d",
+  "webgpu/shader/execution/binary/f32_subtraction.bin": "867a37cd",
+  "webgpu/shader/execution/binary/i32_arithmetic.bin": "300eab87",
+  "webgpu/shader/execution/binary/i32_comparison.bin": "efa097b6",
+  "webgpu/shader/execution/binary/u32_arithmetic.bin": "a83b5a8",
+  "webgpu/shader/execution/binary/u32_comparison.bin": "37aae6eb",
+  "webgpu/shader/execution/abs.bin": "74aa7896",
+  "webgpu/shader/execution/acos.bin": "df364e7",
+  "webgpu/shader/execution/acosh.bin": "218177b5",
+  "webgpu/shader/execution/asin.bin": "cf90fb40",
+  "webgpu/shader/execution/asinh.bin": "e88426c0",
+  "webgpu/shader/execution/atan.bin": "30e5cdba",
+  "webgpu/shader/execution/atan2.bin": "2cb9c01c",
+  "webgpu/shader/execution/atanh.bin": "cf945e3e",
+  "webgpu/shader/execution/bitcast.bin": "3a3d8e53",
+  "webgpu/shader/execution/ceil.bin": "c89008e0",
+  "webgpu/shader/execution/clamp.bin": "cf06c601",
+  "webgpu/shader/execution/cos.bin": "2bbc33",
+  "webgpu/shader/execution/cosh.bin": "6980ff40",
+  "webgpu/shader/execution/cross.bin": "16eba697",
+  "webgpu/shader/execution/degrees.bin": "b7b360b3",
+  "webgpu/shader/execution/determinant.bin": "c35288bd",
+  "webgpu/shader/execution/distance.bin": "a5c99ede",
+  "webgpu/shader/execution/dot.bin": "9ed123dd",
+  "webgpu/shader/execution/exp.bin": "d594d140",
+  "webgpu/shader/execution/exp2.bin": "4b15e525",
+  "webgpu/shader/execution/faceForward.bin": "8c5e4c59",
+  "webgpu/shader/execution/floor.bin": "f2058b59",
+  "webgpu/shader/execution/fma.bin": "295ebf2",
+  "webgpu/shader/execution/fract.bin": "7dc6ac86",
+  "webgpu/shader/execution/frexp.bin": "f33cefc",
+  "webgpu/shader/execution/inverseSqrt.bin": "96a23362",
+  "webgpu/shader/execution/ldexp.bin": "750f20dc",
+  "webgpu/shader/execution/length.bin": "5ff28a30",
+  "webgpu/shader/execution/log.bin": "4a89a92e",
+  "webgpu/shader/execution/log2.bin": "fc5c8930",
+  "webgpu/shader/execution/max.bin": "d5ec21f8",
+  "webgpu/shader/execution/min.bin": "3d162b98",
+  "webgpu/shader/execution/mix.bin": "d265c17b",
+  "webgpu/shader/execution/modf.bin": "bdbf837e",
+  "webgpu/shader/execution/normalize.bin": "1493faef",
+  "webgpu/shader/execution/pack2x16float.bin": "db16cd5e",
+  "webgpu/shader/execution/pow.bin": "e170f03b",
+  "webgpu/shader/execution/quantizeToF16.bin": "5f6adf6a",
+  "webgpu/shader/execution/radians.bin": "ba8aa775",
+  "webgpu/shader/execution/reflect.bin": "5a44a6bc",
+  "webgpu/shader/execution/refract.bin": "20905378",
+  "webgpu/shader/execution/round.bin": "f3ed343",
+  "webgpu/shader/execution/saturate.bin": "83255982",
+  "webgpu/shader/execution/sign.bin": "89153648",
+  "webgpu/shader/execution/sin.bin": "a0ce578b",
+  "webgpu/shader/execution/sinh.bin": "dacde8f8",
+  "webgpu/shader/execution/smoothstep.bin": "497d430",
+  "webgpu/shader/execution/sqrt.bin": "1c00b1c",
+  "webgpu/shader/execution/step.bin": "c4432092",
+  "webgpu/shader/execution/tan.bin": "26d0a994",
+  "webgpu/shader/execution/tanh.bin": "d16dc4ff",
+  "webgpu/shader/execution/transpose.bin": "edbfdf6e",
+  "webgpu/shader/execution/trunc.bin": "1bc2032",
+  "webgpu/shader/execution/unpack2x16float.bin": "c7cbb7f9",
+  "webgpu/shader/execution/unpack2x16snorm.bin": "4673a",
+  "webgpu/shader/execution/unpack2x16unorm.bin": "5db86731",
+  "webgpu/shader/execution/unpack4x8snorm.bin": "a99a3399",
+  "webgpu/shader/execution/unpack4x8unorm.bin": "26831332",
+  "webgpu/shader/execution/unary/af_arithmetic.bin": "e82077ef",
+  "webgpu/shader/execution/unary/af_assignment.bin": "18c7424a",
+  "webgpu/shader/execution/unary/bool_conversion.bin": "4405b038",
+  "webgpu/shader/execution/unary/f16_arithmetic.bin": "717193f7",
+  "webgpu/shader/execution/unary/f16_conversion.bin": "106fc817",
+  "webgpu/shader/execution/unary/f32_arithmetic.bin": "a1123962",
+  "webgpu/shader/execution/unary/f32_conversion.bin": "ab27a511",
+  "webgpu/shader/execution/unary/i32_arithmetic.bin": "5486010b",
+  "webgpu/shader/execution/unary/i32_complement.bin": "ec548b43",
+  "webgpu/shader/execution/unary/i32_conversion.bin": "e77636a7",
+  "webgpu/shader/execution/unary/u32_complement.bin": "64732842",
+  "webgpu/shader/execution/unary/u32_conversion.bin": "9d69dfe1"
 }
\ 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 808ddd88bd19..d019622c3786 100644
Binary files a/src/resources/cache/webgpu/shader/execution/bitcast.bin and b/src/resources/cache/webgpu/shader/execution/bitcast.bin differ
diff --git a/src/resources/four-colors-h264-bt601-hflip.mp4 b/src/resources/four-colors-h264-bt601-hflip.mp4
new file mode 100644
index 000000000000..f83b4f96985f
Binary files /dev/null and b/src/resources/four-colors-h264-bt601-hflip.mp4 differ
diff --git a/src/resources/four-colors-h264-bt601-vflip.mp4 b/src/resources/four-colors-h264-bt601-vflip.mp4
new file mode 100644
index 000000000000..90c3297a9a2e
Binary files /dev/null and b/src/resources/four-colors-h264-bt601-vflip.mp4 differ
diff --git a/src/resources/four-colors-vp9-bt601-hflip.mp4 b/src/resources/four-colors-vp9-bt601-hflip.mp4
new file mode 100644
index 000000000000..f782c326517d
Binary files /dev/null and b/src/resources/four-colors-vp9-bt601-hflip.mp4 differ
diff --git a/src/resources/four-colors-vp9-bt601-vflip.mp4 b/src/resources/four-colors-vp9-bt601-vflip.mp4
new file mode 100644
index 000000000000..c9de14696a73
Binary files /dev/null and b/src/resources/four-colors-vp9-bt601-vflip.mp4 differ
diff --git a/src/webgpu/api/operation/storage_texture/read_write.spec.ts b/src/webgpu/api/operation/storage_texture/read_write.spec.ts
new file mode 100644
index 000000000000..b0a2256d9243
--- /dev/null
+++ b/src/webgpu/api/operation/storage_texture/read_write.spec.ts
@@ -0,0 +1,377 @@
+export const description = `
+Tests for the behavior of read-write storage textures.
+
+TODO:
+- Test resource usage transitions with read-write storage textures
+`;
+
+import { makeTestGroup } from '../../../../common/framework/test_group.js';
+import { assert, unreachable } from '../../../../common/util/util.js';
+import { kTextureDimensions } from '../../../capability_info.js';
+import { kColorTextureFormats, kTextureFormatInfo } from '../../../format_info.js';
+import { GPUTest } from '../../../gpu_test.js';
+import { align } from '../../../util/math.js';
+
+const kShaderStagesForReadWriteStorageTexture = ['fragment', 'compute'] as const;
+type ShaderStageForReadWriteStorageTexture =
+  (typeof kShaderStagesForReadWriteStorageTexture)[number];
+
+class F extends GPUTest {
+  GetInitialData(storageTexture: GPUTexture): ArrayBuffer {
+    const format = storageTexture.format;
+    const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock;
+    assert(bytesPerBlock !== undefined);
+
+    const width = storageTexture.width;
+    const height = storageTexture.height;
+    const depthOrArrayLayers = storageTexture.depthOrArrayLayers;
+    const initialData = new ArrayBuffer(bytesPerBlock * width * height * depthOrArrayLayers);
+    const initialTypedData = this.GetTypedArrayBuffer(initialData, format);
+    for (let z = 0; z < depthOrArrayLayers; ++z) {
+      for (let y = 0; y < height; ++y) {
+        for (let x = 0; x < width; ++x) {
+          const index = z * width * height + y * width + x;
+          switch (format) {
+            case 'r32sint':
+              initialTypedData[index] = (index & 1 ? 1 : -1) * (2 * index + 1);
+              break;
+            case 'r32uint':
+              initialTypedData[index] = 2 * index + 1;
+              break;
+            case 'r32float':
+              initialTypedData[index] = (2 * index + 1) / 10.0;
+              break;
+          }
+        }
+      }
+    }
+    return initialData;
+  }
+
+  GetTypedArrayBuffer(arrayBuffer: ArrayBuffer, format: GPUTextureFormat) {
+    switch (format) {
+      case 'r32sint':
+        return new Int32Array(arrayBuffer);
+      case 'r32uint':
+        return new Uint32Array(arrayBuffer);
+      case 'r32float':
+        return new Float32Array(arrayBuffer);
+      default:
+        unreachable();
+        return new Uint8Array(arrayBuffer);
+    }
+  }
+
+  GetExpectedData(
+    shaderStage: ShaderStageForReadWriteStorageTexture,
+    storageTexture: GPUTexture,
+    initialData: ArrayBuffer
+  ): ArrayBuffer {
+    const format = storageTexture.format;
+    const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock;
+    assert(bytesPerBlock !== undefined);
+
+    const width = storageTexture.width;
+    const height = storageTexture.height;
+    const depthOrArrayLayers = storageTexture.depthOrArrayLayers;
+    const bytesPerRowAlignment = align(bytesPerBlock * width, 256);
+    const itemsPerRow = bytesPerRowAlignment / bytesPerBlock;
+
+    const expectedData = new ArrayBuffer(
+      bytesPerRowAlignment * (height * depthOrArrayLayers - 1) + bytesPerBlock * width
+    );
+    const expectedTypedData = this.GetTypedArrayBuffer(expectedData, format);
+    const initialTypedData = this.GetTypedArrayBuffer(initialData, format);
+    for (let z = 0; z < depthOrArrayLayers; ++z) {
+      for (let y = 0; y < height; ++y) {
+        for (let x = 0; x < width; ++x) {
+          const expectedIndex = z * itemsPerRow * height + y * itemsPerRow + x;
+          switch (shaderStage) {
+            case 'compute': {
+              // In the compute shader we flip the texture along the diagonal.
+              const initialIndex =
+                (depthOrArrayLayers - 1 - z) * width * height +
+                (height - 1 - y) * width +
+                (width - 1 - x);
+              expectedTypedData[expectedIndex] = initialTypedData[initialIndex];
+              break;
+            }
+            case 'fragment': {
+              // In the fragment shader we double the original texel value of the read-write storage
+              // texture.
+              const initialIndex = z * width * height + y * width + x;
+              expectedTypedData[expectedIndex] = initialTypedData[initialIndex] * 2;
+              break;
+            }
+          }
+        }
+      }
+    }
+    return expectedData;
+  }
+
+  RecordCommandsToTransform(
+    device: GPUDevice,
+    shaderStage: ShaderStageForReadWriteStorageTexture,
+    commandEncoder: GPUCommandEncoder,
+    rwTexture: GPUTexture
+  ) {
+    let declaration = '';
+    switch (rwTexture.dimension) {
+      case '1d':
+        declaration = 'texture_storage_1d';
+        break;
+      case '2d':
+        declaration =
+          rwTexture.depthOrArrayLayers > 1 ? 'texture_storage_2d_array' : 'texture_storage_2d';
+        break;
+      case '3d':
+        declaration = 'texture_storage_3d';
+        break;
+    }
+    const textureDeclaration = `
+    @group(0) @binding(0) var rwTexture: ${declaration}<${rwTexture.format}, read_write>;
+    `;
+
+    switch (shaderStage) {
+      case 'fragment': {
+        const vertexShader = `
+        @vertex
+        fn main(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f {
+            var pos = array(
+              vec2f(-1.0, -1.0),
+              vec2f(-1.0,  1.0),
+              vec2f( 1.0, -1.0),
+              vec2f(-1.0,  1.0),
+              vec2f( 1.0, -1.0),
+              vec2f( 1.0,  1.0));
+            return vec4f(pos[VertexIndex], 0.0, 1.0);
+        }
+        `;
+        let textureLoadStoreCoord = '';
+        switch (rwTexture.dimension) {
+          case '1d':
+            textureLoadStoreCoord = 'textureCoord.x';
+            break;
+          case '2d':
+            textureLoadStoreCoord =
+              rwTexture.depthOrArrayLayers > 1 ? 'textureCoord, z' : 'textureCoord';
+            break;
+          case '3d':
+            textureLoadStoreCoord = 'vec3u(textureCoord, z)';
+            break;
+        }
+        const fragmentShader = `
+        ${textureDeclaration}
+        @fragment
+        fn main(@builtin(position) fragCoord: vec4f) -> @location(0) vec4f {
+          let textureCoord = vec2u(fragCoord.xy);
+
+          for (var z = 0u; z < ${rwTexture.depthOrArrayLayers}; z++) {
+            let initialValue = textureLoad(rwTexture, ${textureLoadStoreCoord});
+            let outputValue = initialValue * 2;
+            textureStore(rwTexture, ${textureLoadStoreCoord}, outputValue);
+          }
+
+          return vec4f(0.0, 1.0, 0.0, 1.0);
+        }
+        `;
+        const renderPipeline = device.createRenderPipeline({
+          layout: 'auto',
+          vertex: {
+            module: device.createShaderModule({
+              code: vertexShader,
+            }),
+          },
+          fragment: {
+            module: device.createShaderModule({
+              code: fragmentShader,
+            }),
+            targets: [
+              {
+                format: 'rgba8unorm',
+              },
+            ],
+          },
+          primitive: {
+            topology: 'triangle-list',
+          },
+        });
+
+        const bindGroup = device.createBindGroup({
+          layout: renderPipeline.getBindGroupLayout(0),
+          entries: [
+            {
+              binding: 0,
+              resource: rwTexture.createView(),
+            },
+          ],
+        });
+
+        const dummyColorTexture = device.createTexture({
+          size: [rwTexture.width, rwTexture.height, 1],
+          usage: GPUTextureUsage.RENDER_ATTACHMENT,
+          format: 'rgba8unorm',
+        });
+
+        const renderPassEncoder = commandEncoder.beginRenderPass({
+          colorAttachments: [
+            {
+              view: dummyColorTexture.createView(),
+              loadOp: 'clear',
+              clearValue: { r: 0, g: 0, b: 0, a: 0 },
+              storeOp: 'store',
+            },
+          ],
+        });
+        renderPassEncoder.setPipeline(renderPipeline);
+        renderPassEncoder.setBindGroup(0, bindGroup);
+        renderPassEncoder.draw(6);
+        renderPassEncoder.end();
+        break;
+      }
+      case 'compute': {
+        let textureLoadCoord = '';
+        let textureStoreCoord = '';
+        switch (rwTexture.dimension) {
+          case '1d':
+            textureLoadCoord = 'dimension - 1u - invocationID.x';
+            textureStoreCoord = 'invocationID.x';
+            break;
+          case '2d':
+            textureLoadCoord =
+              rwTexture.depthOrArrayLayers > 1
+                ? `vec2u(dimension.x - 1u - invocationID.x, dimension.y - 1u - invocationID.y),
+                   textureNumLayers(rwTexture) - 1u - invocationID.z`
+                : `vec2u(dimension.x - 1u - invocationID.x, dimension.y - 1u - invocationID.y)`;
+            textureStoreCoord =
+              rwTexture.depthOrArrayLayers > 1
+                ? 'invocationID.xy, invocationID.z'
+                : 'invocationID.xy';
+            break;
+          case '3d':
+            textureLoadCoord = `
+              vec3u(dimension.x - 1u - invocationID.x, dimension.y - 1u - invocationID.y,
+                    dimension.z - 1u - invocationID.z)`;
+            textureStoreCoord = 'invocationID';
+            break;
+        }
+
+        const computeShader = `
+        ${textureDeclaration}
+        @compute
+        @workgroup_size(${rwTexture.width}, ${rwTexture.height}, ${rwTexture.depthOrArrayLayers})
+        fn main(@builtin(local_invocation_id) invocationID: vec3u) {
+          let dimension = textureDimensions(rwTexture);
+
+          let initialValue = textureLoad(rwTexture, ${textureLoadCoord});
+          textureBarrier();
+
+          textureStore(rwTexture, ${textureStoreCoord}, initialValue);
+        }`;
+
+        const computePipeline = device.createComputePipeline({
+          compute: {
+            module: device.createShaderModule({
+              code: computeShader,
+            }),
+          },
+          layout: 'auto',
+        });
+        const bindGroup = device.createBindGroup({
+          layout: computePipeline.getBindGroupLayout(0),
+          entries: [
+            {
+              binding: 0,
+              resource: rwTexture.createView(),
+            },
+          ],
+        });
+        const computePassEncoder = commandEncoder.beginComputePass();
+        computePassEncoder.setPipeline(computePipeline);
+        computePassEncoder.setBindGroup(0, bindGroup);
+        computePassEncoder.dispatchWorkgroups(1);
+        computePassEncoder.end();
+        break;
+      }
+    }
+  }
+}
+
+export const g = makeTestGroup(F);
+
+g.test('basic')
+  .desc(
+    `The basic functionality tests for read-write storage textures. In the test we read data from
+    the read-write storage texture, do transforms and write the data back to the read-write storage
+    texture. textureBarrier() is also called in the tests using compute pipelines.`
+  )
+  .params(u =>
+    u
+      .combine('format', kColorTextureFormats)
+      .filter(p => kTextureFormatInfo[p.format].color?.readWriteStorage === true)
+      .combine('shaderStage', kShaderStagesForReadWriteStorageTexture)
+      .combine('textureDimension', kTextureDimensions)
+      .combine('depthOrArrayLayers', [1, 2] as const)
+      .unless(p => p.textureDimension === '1d' && p.depthOrArrayLayers > 1)
+  )
+  .fn(t => {
+    const { format, shaderStage, textureDimension, depthOrArrayLayers } = t.params;
+
+    const kWidth = 16;
+    const height = textureDimension === '1d' ? 1 : 8;
+    const textureSize = [kWidth, height, depthOrArrayLayers] as const;
+    const storageTexture = t.device.createTexture({
+      format,
+      dimension: textureDimension,
+      size: textureSize,
+      usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING,
+    });
+
+    const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock;
+    const initialData = t.GetInitialData(storageTexture);
+    t.queue.writeTexture(
+      { texture: storageTexture },
+      initialData,
+      {
+        bytesPerRow: bytesPerBlock * kWidth,
+        rowsPerImage: height,
+      },
+      textureSize
+    );
+
+    const commandEncoder = t.device.createCommandEncoder();
+
+    t.RecordCommandsToTransform(t.device, shaderStage, commandEncoder, storageTexture);
+
+    const expectedData = t.GetExpectedData(shaderStage, storageTexture, initialData);
+    const readbackBuffer = t.device.createBuffer({
+      size: expectedData.byteLength,
+      usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
+    });
+    const bytesPerRow = align(bytesPerBlock * kWidth, 256);
+    commandEncoder.copyTextureToBuffer(
+      {
+        texture: storageTexture,
+      },
+      {
+        buffer: readbackBuffer,
+        bytesPerRow,
+        rowsPerImage: height,
+      },
+      textureSize
+    );
+    t.queue.submit([commandEncoder.finish()]);
+
+    switch (format) {
+      case 'r32sint':
+        t.expectGPUBufferValuesEqual(readbackBuffer, new Int32Array(expectedData));
+        break;
+      case 'r32uint':
+        t.expectGPUBufferValuesEqual(readbackBuffer, new Uint32Array(expectedData));
+        break;
+      case 'r32float':
+        t.expectGPUBufferValuesEqual(readbackBuffer, new Float32Array(expectedData));
+        break;
+    }
+  });
diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts
index 1bd5d3b7c698..d26d93ca2e8d 100644
--- a/src/webgpu/capability_info.ts
+++ b/src/webgpu/capability_info.ts
@@ -821,6 +821,7 @@ export const kKnownWGSLLanguageFeatures = [
   'readonly_and_readwrite_storage_textures',
   'packed_4x8_integer_dot_product',
   'unrestricted_pointer_parameters',
+  'pointer_composite_access',
 ] as const;
 
 export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number];
diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index 933ed575565c..cf8acba865de 100644
--- a/src/webgpu/listing_meta.json
+++ b/src/webgpu/listing_meta.json
@@ -199,6 +199,7 @@
   "webgpu:api,operation,shader_module,compilation_info:getCompilationInfo_returns:*": { "subcaseMS": 0.284 },
   "webgpu:api,operation,shader_module,compilation_info:line_number_and_position:*": { "subcaseMS": 1.867 },
   "webgpu:api,operation,shader_module,compilation_info:offset_and_length:*": { "subcaseMS": 1.648 },
+  "webgpu:api,operation,storage_texture,read_write:basic:*": { "subcaseMS": 5.000 },
   "webgpu:api,operation,texture_view,format_reinterpretation:render_and_resolve_attachment:*": { "subcaseMS": 14.488 },
   "webgpu:api,operation,texture_view,format_reinterpretation:texture_binding:*": { "subcaseMS": 17.225 },
   "webgpu:api,operation,texture_view,read:aspect:*": { "subcaseMS": 0.601 },
@@ -1226,6 +1227,8 @@
   "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 },
@@ -1507,6 +1510,13 @@
   "webgpu:shader,execution,expression,call,builtin,unpack4x8unorm:unpack:*": { "subcaseMS": 11.776 },
   "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: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: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,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 },
@@ -1771,22 +1781,22 @@
   "webgpu:shader,validation,expression,call,builtin,log:values:*": { "subcaseMS": 0.291 },
   "webgpu:shader,validation,expression,call,builtin,modf:integer_argument:*": { "subcaseMS": 1.089 },
   "webgpu:shader,validation,expression,call,builtin,modf:values:*": { "subcaseMS": 1.866 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8:bad_args:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8:must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8:supported:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8:unsupported:*": { "subcaseMS": 0.300 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:bad_args:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:supported:*": { "subcaseMS": 0.101 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:unsupported:*": { "subcaseMS": 0.300 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8:bad_args:*": { "subcaseMS": 0.200 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8:must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8:supported:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8:unsupported:*": { "subcaseMS": 0.300 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:bad_args:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:supported:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:unsupported:*": { "subcaseMS": 0.351 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8:bad_args:*": { "subcaseMS": 40.750 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8:must_use:*": { "subcaseMS": 6.500 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8:supported:*": { "subcaseMS": 113.501 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8:unsupported:*": { "subcaseMS": 739.400 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:bad_args:*": { "subcaseMS": 39.240 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:must_use:*": { "subcaseMS": 34.301 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:supported:*": { "subcaseMS": 100.450 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xI8Clamp:unsupported:*": { "subcaseMS": 751.101 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8:bad_args:*": { "subcaseMS": 37.770 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8:must_use:*": { "subcaseMS": 5.300 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8:supported:*": { "subcaseMS": 449.800 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8:unsupported:*": { "subcaseMS": 773.702 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:bad_args:*": { "subcaseMS": 124.860 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:must_use:*": { "subcaseMS": 32.600 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:supported:*": { "subcaseMS": 134.750 },
+  "webgpu:shader,validation,expression,call,builtin,pack4xU8Clamp:unsupported:*": { "subcaseMS": 570.500 },
   "webgpu:shader,validation,expression,call,builtin,radians:integer_argument:*": { "subcaseMS": 1.811 },
   "webgpu:shader,validation,expression,call,builtin,radians:values:*": { "subcaseMS": 0.382 },
   "webgpu:shader,validation,expression,call,builtin,round:integer_argument:*": { "subcaseMS": 1.834 },
@@ -1803,26 +1813,23 @@
   "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,unpack4xI8:bad_args:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xI8:must_use:*": { "subcaseMS": 0.000 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xI8:supported:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xI8:unsupported:*": { "subcaseMS": 0.351 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xU8:bad_args:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xU8:must_use:*": { "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: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,validation,expression,call,builtin,unpack4xU8:supported:*": { "subcaseMS": 0.100 },
-  "webgpu:shader,validation,expression,call,builtin,unpack4xU8:unsupported:*": { "subcaseMS": 0.301 },
+  "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 },
+  "webgpu:shader,validation,expression,call,builtin,unpack4xI8:unsupported:*": { "subcaseMS": 615.301 },
+  "webgpu:shader,validation,expression,call,builtin,unpack4xU8:bad_args:*": { "subcaseMS": 40.113 },
+  "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,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_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_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_pointers:*": { "subcaseMS": 1.537 },
   "webgpu:shader,validation,functions,restrictions:call_arg_types_match_params:*": { "subcaseMS": 1.518 },
   "webgpu:shader,validation,functions,restrictions:entry_point_call_target:*": { "subcaseMS": 1.734 },
diff --git a/src/webgpu/shader/execution/expression/call/builtin/dot4I8Packed.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/dot4I8Packed.spec.ts
new file mode 100644
index 000000000000..dee5290281a8
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/dot4I8Packed.spec.ts
@@ -0,0 +1,74 @@
+export const description = `
+Execution tests for the 'dot4I8Packed' builtin function
+
+@const fn dot4I8Packed(e1: u32 ,e2: u32) -> i32
+e1 and e2 are interpreted as vectors with four 8-bit signed integer components. Return the signed
+integer dot product of these two vectors. Each component is sign-extended to i32 before performing
+the multiply, and then the add operations are done in WGSL i32 with wrapping behaviour.
+`;
+
+import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { GPUTest } from '../../../../../gpu_test.js';
+import { TypeI32, TypeU32, i32, u32 } from '../../../../../util/conversion.js';
+import { Case } from '../../case.js';
+import { allInputSources, Config, run } from '../../expression.js';
+
+import { builtin } from './builtin.js';
+
+export const g = makeTestGroup(GPUTest);
+
+g.test('basic')
+  .specURL('https://www.w3.org/TR/WGSL/#dot4I8Packed-builtin')
+  .desc(
+    `
+@const fn dot4I8Packed(e1: u32, e2: u32) -> i32
+  `
+  )
+  .params(u => u.combine('inputSource', allInputSources))
+  .fn(async t => {
+    const cfg: Config = t.params;
+
+    const dot4I8Packed = (e1: number, e2: number) => {
+      let result = 0;
+      for (let i = 0; i < 4; ++i) {
+        let e1_i = (e1 >> (i * 8)) & 0xff;
+        if (e1_i >= 128) {
+          e1_i -= 256;
+        }
+        let e2_i = (e2 >> (i * 8)) & 0xff;
+        if (e2_i >= 128) {
+          e2_i -= 256;
+        }
+        result += e1_i * e2_i;
+      }
+      return result;
+    };
+
+    const testInputs = [
+      // dot({0, 0, 0, 0}, {0, 0, 0, 0})
+      [0, 0],
+      // dot({127, 127, 127, 127}, {127, 127, 127, 127})
+      [0x7f7f7f7f, 0x7f7f7f7f],
+      // dot({-128, -128, -128, -128}, {-128, -128, -128, -128})
+      [0x80808080, 0x80808080],
+      // dot({127, 127, 127, 127}, {-128, -128, -128, -128})
+      [0x7f7f7f7f, 0x80808080],
+      // dot({1, 2, 3, 4}, {5, 6, 7, 8})
+      [0x01020304, 0x05060708],
+      // dot({1, 2, 3, 4}, {-1, -2, -3, -4})
+      [0x01020304, 0xfffefdfc],
+      // dot({-5, -6, -7, -8}, {5, 6, 7, 8})
+      [0xfbfaf9f8, 0x05060708],
+      // dot({-9, -10, -11, -12}, {-13, -14, -15, -16})
+      [0xf7f6f5f4, 0xf3f2f1f0],
+    ] as const;
+
+    const makeCase = (x: number, y: number): Case => {
+      return { input: [u32(x), u32(y)], expected: i32(dot4I8Packed(x, y)) };
+    };
+    const cases: Array<Case> = testInputs.flatMap(v => {
+      return [makeCase(...(v as [number, number]))];
+    });
+
+    await run(t, builtin('dot4I8Packed'), [TypeU32, TypeU32], TypeI32, cfg, cases);
+  });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/dot4U8Packed.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/dot4U8Packed.spec.ts
new file mode 100644
index 000000000000..f0dd6fc5081b
--- /dev/null
+++ b/src/webgpu/shader/execution/expression/call/builtin/dot4U8Packed.spec.ts
@@ -0,0 +1,59 @@
+export const description = `
+Execution tests for the 'dot4U8Packed' builtin function
+
+@const fn dot4U8Packed(e1: u32 ,e2: u32) -> u32
+e1 and e2 are interpreted as vectors with four 8-bit unsigned integer components. Return the
+unsigned integer dot product of these two vectors.
+`;
+
+import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { GPUTest } from '../../../../../gpu_test.js';
+import { TypeU32, u32 } from '../../../../../util/conversion.js';
+import { Case } from '../../case.js';
+import { allInputSources, Config, run } from '../../expression.js';
+
+import { builtin } from './builtin.js';
+
+export const g = makeTestGroup(GPUTest);
+
+g.test('basic')
+  .specURL('https://www.w3.org/TR/WGSL/#dot4U8Packed-builtin')
+  .desc(
+    `
+@const fn dot4U8Packed(e1: u32, e2: u32) -> u32
+  `
+  )
+  .params(u => u.combine('inputSource', allInputSources))
+  .fn(async t => {
+    const cfg: Config = t.params;
+
+    const dot4U8Packed = (e1: number, e2: number) => {
+      let result = 0;
+      for (let i = 0; i < 4; ++i) {
+        const e1_i = (e1 >> (i * 8)) & 0xff;
+        const e2_i = (e2 >> (i * 8)) & 0xff;
+        result += e1_i * e2_i;
+      }
+      return result;
+    };
+
+    const testInputs = [
+      // dot({0, 0, 0, 0}, {0, 0, 0, 0})
+      [0, 0],
+      // dot({255u, 255u, 255u, 255u}, {255u, 255u, 255u, 255u})
+      [0xffffffff, 0xffffffff],
+      // dot({1u, 2u, 3u, 4u}, {5u, 6u, 7u, 8u})
+      [0x01020304, 0x05060708],
+      // dot({120u, 90u, 60u, 30u}, {50u, 100u, 150u, 200u})
+      [0x785a3c1e, 0x326496c8],
+    ] as const;
+
+    const makeCase = (x: number, y: number): Case => {
+      return { input: [u32(x), u32(y)], expected: u32(dot4U8Packed(x, y)) };
+    };
+    const cases: Array<Case> = testInputs.flatMap(v => {
+      return [makeCase(...(v as [number, number]))];
+    });
+
+    await run(t, builtin('dot4U8Packed'), [TypeU32, TypeU32], TypeU32, cfg, cases);
+  });
diff --git a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts
index fcf3159c642c..a40b42633283 100644
--- a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts
+++ b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts
@@ -1,7 +1,6 @@
 export const description = `Test compute shader builtin variables`;
 
 import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { iterRange } from '../../../../common/util/util.js';
 import { GPUTest } from '../../../gpu_test.js';
 
 export const g = makeTestGroup(GPUTest);
@@ -98,17 +97,14 @@ g.test('inputs')
 
     // WGSL shader that stores every builtin value to a buffer, for every invocation in the grid.
     const wgsl = `
-      struct S {
-        data : array<u32>
+      struct Outputs {
+        local_id: vec3u,
+        local_index: u32,
+        global_id: vec3u,
+        group_id: vec3u,
+        num_groups: vec3u,
       };
-      struct V {
-        data : array<vec3<u32>>
-      };
-      @group(0) @binding(0) var<storage, read_write> local_id_out : V;
-      @group(0) @binding(1) var<storage, read_write> local_index_out : S;
-      @group(0) @binding(2) var<storage, read_write> global_id_out : V;
-      @group(0) @binding(3) var<storage, read_write> group_id_out : V;
-      @group(0) @binding(4) var<storage, read_write> num_groups_out : V;
+      @group(0) @binding(0) var<storage, read_write> outputs : array<Outputs>;
 
       ${structures}
 
@@ -122,11 +118,13 @@ g.test('inputs')
         ) {
         let group_index = ((${group_id}.z * ${num_groups}.y) + ${group_id}.y) * ${num_groups}.x + ${group_id}.x;
         let global_index = group_index * ${invocationsPerGroup}u + ${local_index};
-        local_id_out.data[global_index] = ${local_id};
-        local_index_out.data[global_index] = ${local_index};
-        global_id_out.data[global_index] = ${global_id};
-        group_id_out.data[global_index] = ${group_id};
-        num_groups_out.data[global_index] = ${num_groups};
+        var o: Outputs;
+        o.local_id = ${local_id};
+        o.local_index = ${local_index};
+        o.global_id = ${global_id};
+        o.group_id = ${group_id};
+        o.num_groups = ${num_groups};
+        outputs[global_index] = o;
       }
     `;
 
@@ -140,35 +138,24 @@ g.test('inputs')
       },
     });
 
-    // Helper to create a `size`-byte buffer with binding number `binding`.
-    function createBuffer(size: number, binding: number) {
-      const buffer = t.device.createBuffer({
-        size,
-        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
-      });
-      t.trackForCleanup(buffer);
-
-      bindGroupEntries.push({
-        binding,
-        resource: {
-          buffer,
-        },
-      });
-
-      return buffer;
-    }
+    // Offsets are in u32 size units
+    const kLocalIdOffset = 0;
+    const kLocalIndexOffset = 3;
+    const kGlobalIdOffset = 4;
+    const kGroupIdOffset = 8;
+    const kNumGroupsOffset = 12;
+    const kOutputElementSize = 16;
 
     // Create the output buffers.
-    const bindGroupEntries: GPUBindGroupEntry[] = [];
-    const localIdBuffer = createBuffer(totalInvocations * 16, 0);
-    const localIndexBuffer = createBuffer(totalInvocations * 4, 1);
-    const globalIdBuffer = createBuffer(totalInvocations * 16, 2);
-    const groupIdBuffer = createBuffer(totalInvocations * 16, 3);
-    const numGroupsBuffer = createBuffer(totalInvocations * 16, 4);
+    const outputBuffer = t.device.createBuffer({
+      size: totalInvocations * kOutputElementSize * 4,
+      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
+    });
+    t.trackForCleanup(outputBuffer);
 
     const bindGroup = t.device.createBindGroup({
       layout: pipeline.getBindGroupLayout(0),
-      entries: bindGroupEntries,
+      entries: [{ binding: 0, resource: { buffer: outputBuffer } }],
     });
 
     // Run the shader.
@@ -204,11 +191,7 @@ g.test('inputs')
     // Helper to check that the vec3<u32> value at each index of the provided `output` buffer
     // matches the expected value for that invocation, as generated by the `getBuiltinValue`
     // function. The `name` parameter is the builtin name, used for error messages.
-    const checkEachIndex = (
-      output: Uint32Array,
-      name: string,
-      getBuiltinValue: (groupId: vec3, localId: vec3) => vec3
-    ) => {
+    const checkEachIndex = (output: Uint32Array) => {
       // Loop over workgroups.
       for (let gz = 0; gz < t.params.numGroups.z; gz++) {
         for (let gy = 0; gy < t.params.numGroups.y; gy++) {
@@ -220,30 +203,44 @@ g.test('inputs')
                   const groupIndex = (gz * t.params.numGroups.y + gy) * t.params.numGroups.x + gx;
                   const localIndex = (lz * t.params.groupSize.y + ly) * t.params.groupSize.x + lx;
                   const globalIndex = groupIndex * invocationsPerGroup + localIndex;
-                  const expected = getBuiltinValue(
-                    { x: gx, y: gy, z: gz },
-                    { x: lx, y: ly, z: lz }
-                  );
-                  if (output[globalIndex * 4 + 0] !== expected.x) {
-                    return new Error(
-                      `${name}.x failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` +
-                        `    expected: ${expected.x}\n` +
-                        `    got:      ${output[globalIndex * 4 + 0]}`
-                    );
-                  }
-                  if (output[globalIndex * 4 + 1] !== expected.y) {
-                    return new Error(
-                      `${name}.y failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` +
-                        `    expected: ${expected.y}\n` +
-                        `    got:      ${output[globalIndex * 4 + 1]}`
+                  const globalOffset = globalIndex * kOutputElementSize;
+
+                  const expectEqual = (name: string, expected: number, actual: number) => {
+                    if (actual !== expected) {
+                      return new Error(
+                        `${name} failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` +
+                          `    expected: ${expected}\n` +
+                          `    got:      ${actual}`
+                      );
+                    }
+                    return undefined;
+                  };
+
+                  const checkVec3Value = (name: string, fieldOffset: number, expected: vec3) => {
+                    const offset = globalOffset + fieldOffset;
+                    return (
+                      expectEqual(`${name}.x`, expected.x, output[offset + 0]) ||
+                      expectEqual(`${name}.y`, expected.y, output[offset + 1]) ||
+                      expectEqual(`${name}.z`, expected.z, output[offset + 2])
                     );
-                  }
-                  if (output[globalIndex * 4 + 2] !== expected.z) {
-                    return new Error(
-                      `${name}.z failed at group(${gx},${gy},${gz}) local(${lx},${ly},${lz}))\n` +
-                        `    expected: ${expected.z}\n` +
-                        `    got:      ${output[globalIndex * 4 + 2]}`
+                  };
+
+                  const error =
+                    checkVec3Value('local_id', kLocalIdOffset, { x: lx, y: ly, z: lz }) ||
+                    checkVec3Value('global_id', kGlobalIdOffset, {
+                      x: gx * t.params.groupSize.x + lx,
+                      y: gy * t.params.groupSize.y + ly,
+                      z: gz * t.params.groupSize.z + lz,
+                    }) ||
+                    checkVec3Value('group_id', kGroupIdOffset, { x: gx, y: gy, z: gz }) ||
+                    checkVec3Value('num_groups', kNumGroupsOffset, t.params.numGroups) ||
+                    expectEqual(
+                      'local_index',
+                      localIndex,
+                      output[globalOffset + kLocalIndexOffset]
                     );
+                  if (error) {
+                    return error;
                   }
                 }
               }
@@ -254,44 +251,8 @@ g.test('inputs')
       return undefined;
     };
 
-    // Check @builtin(local_invocation_index) values.
-    t.expectGPUBufferValuesEqual(
-      localIndexBuffer,
-      new Uint32Array([...iterRange(totalInvocations, x => x % invocationsPerGroup)])
-    );
-
-    // Check @builtin(local_invocation_id) values.
-    t.expectGPUBufferValuesPassCheck(
-      localIdBuffer,
-      outputData => checkEachIndex(outputData, 'local_invocation_id', (_, localId) => localId),
-      { type: Uint32Array, typedLength: totalInvocations * 4 }
-    );
-
-    // Check @builtin(global_invocation_id) values.
-    const getGlobalId = (groupId: vec3, localId: vec3) => {
-      return {
-        x: groupId.x * t.params.groupSize.x + localId.x,
-        y: groupId.y * t.params.groupSize.y + localId.y,
-        z: groupId.z * t.params.groupSize.z + localId.z,
-      };
-    };
-    t.expectGPUBufferValuesPassCheck(
-      globalIdBuffer,
-      outputData => checkEachIndex(outputData, 'global_invocation_id', getGlobalId),
-      { type: Uint32Array, typedLength: totalInvocations * 4 }
-    );
-
-    // Check @builtin(workgroup_id) values.
-    t.expectGPUBufferValuesPassCheck(
-      groupIdBuffer,
-      outputData => checkEachIndex(outputData, 'workgroup_id', (groupId, _) => groupId),
-      { type: Uint32Array, typedLength: totalInvocations * 4 }
-    );
-
-    // Check @builtin(num_workgroups) values.
-    t.expectGPUBufferValuesPassCheck(
-      numGroupsBuffer,
-      outputData => checkEachIndex(outputData, 'num_workgroups', () => t.params.numGroups),
-      { type: Uint32Array, typedLength: totalInvocations * 4 }
-    );
+    t.expectGPUBufferValuesPassCheck(outputBuffer, outputData => checkEachIndex(outputData), {
+      type: Uint32Array,
+      typedLength: outputBuffer.size / 4,
+    });
   });
diff --git a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts
index 0967cf63d8e4..fe74d7efc2e5 100644
--- a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts
+++ b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts
@@ -794,7 +794,7 @@ g.test('inputs,position')
         sampleCount,
         actual,
         expected,
-        maxDiffULPsForFloatFormat: 2,
+        maxDiffULPsForFloatFormat: 4,
       })
     );
   });
@@ -873,7 +873,7 @@ g.test('inputs,interStage')
         sampleCount,
         actual,
         expected,
-        maxDiffULPsForFloatFormat: 3,
+        maxDiffULPsForFloatFormat: 4,
       })
     );
   });
diff --git a/src/webgpu/shader/validation/functions/alias_analysis.spec.ts b/src/webgpu/shader/validation/functions/alias_analysis.spec.ts
index ba39485449df..6b267b67b78b 100644
--- a/src/webgpu/shader/validation/functions/alias_analysis.spec.ts
+++ b/src/webgpu/shader/validation/functions/alias_analysis.spec.ts
@@ -38,50 +38,284 @@ function shouldPass(aliased: boolean, ...uses: UseName[]): boolean {
   return !aliased || !uses.some(u => kUses[u].is_write) || uses.includes('no_access');
 }
 
+type AddressSpace = 'private' | 'function' | 'storage' | 'uniform' | 'workgroup';
+
+const kWritableAddressSpaces = ['private', 'function', 'storage', 'workgroup'] as const;
+
+function ptr(addressSpace: AddressSpace, type: string) {
+  switch (addressSpace) {
+    case 'function':
+      return `ptr<function, ${type}>`;
+    case 'private':
+      return `ptr<private, ${type}>`;
+    case 'storage':
+      return `ptr<storage, ${type}, read_write>`;
+    case 'uniform':
+      return `ptr<uniform, ${type}>`;
+    case 'workgroup':
+      return `ptr<workgroup, ${type}>`;
+  }
+}
+
+function declareModuleScopeVar(
+  name: string,
+  addressSpace: 'private' | 'storage' | 'uniform' | 'workgroup',
+  type: string
+) {
+  const binding = name === 'x' ? 0 : 1;
+  switch (addressSpace) {
+    case 'private':
+      return `var<private> ${name} : ${type};`;
+    case 'storage':
+      return `@binding(${binding}) @group(0) var<storage, read_write> ${name} : ${type};`;
+    case 'uniform':
+      return `@binding(${binding}) @group(0) var<uniform> ${name} : ${type};`;
+    case 'workgroup':
+      return `var<workgroup> ${name} : ${type};`;
+  }
+}
+
+function maybeDeclareModuleScopeVar(name: string, addressSpace: AddressSpace, type: string) {
+  if (addressSpace === 'function') {
+    return '';
+  }
+  return declareModuleScopeVar(name, addressSpace, type);
+}
+
+function maybeDeclareFunctionScopeVar(name: string, addressSpace: AddressSpace, type: string) {
+  switch (addressSpace) {
+    case 'function':
+      return `var ${name} : ${type};`;
+    default:
+      return ``;
+  }
+}
+
+/**
+ * @returns true if a pointer of the given address space requires the
+ * 'unrestricted_pointer_parameters' language feature.
+ */
+function requiresUnrestrictedPointerParameters(addressSpace: AddressSpace) {
+  return addressSpace !== 'function' && addressSpace !== 'private';
+}
+
 g.test('two_pointers')
   .desc(`Test aliasing of two pointers passed to a function.`)
   .params(u =>
     u
-      .combine('address_space', ['private', 'function'] as const)
+      .combine('address_space', kWritableAddressSpaces)
       .combine('a_use', keysOf(kUses))
       .combine('b_use', keysOf(kUses))
       .combine('aliased', [true, false])
       .beginSubcases()
   )
   .fn(t => {
+    if (requiresUnrestrictedPointerParameters(t.params.address_space)) {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+
     const code = `
-${t.params.address_space === 'private' ? `var<private> x : i32; var<private> y : i32;` : ``}
+${maybeDeclareModuleScopeVar('x', t.params.address_space, 'i32')}
+${maybeDeclareModuleScopeVar('y', t.params.address_space, 'i32')}
 
-fn callee(pa : ptr<${t.params.address_space}, i32>,
-          pb : ptr<${t.params.address_space}, i32>) -> i32 {
+fn callee(pa : ${ptr(t.params.address_space, 'i32')},
+          pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
   ${kUses[t.params.a_use].gen(`*pa`)}
   ${kUses[t.params.b_use].gen(`*pb`)}
   return 0;
 }
 
 fn caller() {
-  ${t.params.address_space === 'function' ? `var x : i32; var y : i32;` : ``}
+  ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'i32')}
+  ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'i32')}
   callee(&x, ${t.params.aliased ? `&x` : `&y`});
 }
 `;
     t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code);
   });
 
+g.test('two_pointers_to_array_elements')
+  .desc(`Test aliasing of two array element pointers passed to a function.`)
+  .params(u =>
+    u
+      .combine('address_space', kWritableAddressSpaces)
+      .combine('a_use', keysOf(kUses))
+      .combine('b_use', keysOf(kUses))
+      .combine('index', [0, 1])
+      .combine('aliased', [true, false])
+      .beginSubcases()
+  )
+  .fn(t => {
+    t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+
+    const code = `
+${maybeDeclareModuleScopeVar('x', t.params.address_space, 'array<i32, 4>')}
+${maybeDeclareModuleScopeVar('y', t.params.address_space, 'array<i32, 4>')}
+
+fn callee(pa : ${ptr(t.params.address_space, 'i32')},
+          pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
+  ${kUses[t.params.a_use].gen(`*pa`)}
+  ${kUses[t.params.b_use].gen(`*pb`)}
+  return 0;
+}
+
+fn caller() {
+  ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'array<i32, 4>')}
+  ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'array<i32, 4>')}
+  callee(&x[${t.params.index}], ${t.params.aliased ? `&x[0]` : `&y[0]`});
+}
+`;
+    t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code);
+  });
+
+g.test('two_pointers_to_array_elements_indirect')
+  .desc(
+    `Test aliasing of two array pointers passed to a function, which indexes those arrays and then
+passes the element pointers to another function.`
+  )
+  .params(u =>
+    u
+      .combine('address_space', kWritableAddressSpaces)
+      .combine('a_use', keysOf(kUses))
+      .combine('b_use', keysOf(kUses))
+      .combine('index', [0, 1])
+      .combine('aliased', [true, false])
+      .beginSubcases()
+  )
+  .fn(t => {
+    t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+
+    const code = `
+${maybeDeclareModuleScopeVar('x', t.params.address_space, 'array<i32, 4>')}
+${maybeDeclareModuleScopeVar('y', t.params.address_space, 'array<i32, 4>')}
+
+fn callee(pa : ${ptr(t.params.address_space, 'i32')},
+          pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
+  ${kUses[t.params.a_use].gen(`*pa`)}
+  ${kUses[t.params.b_use].gen(`*pb`)}
+  return 0;
+}
+
+fn index(pa : ${ptr(t.params.address_space, 'array<i32, 4>')},
+         pb : ${ptr(t.params.address_space, 'array<i32, 4>')}) -> i32 {
+  return callee(&(*pa)[${t.params.index}], &(*pb)[0]);
+}
+
+fn caller() {
+  ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'array<i32, 4>')}
+  ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'array<i32, 4>')}
+  index(&x, ${t.params.aliased ? `&x` : `&y`});
+}
+`;
+    t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code);
+  });
+
+g.test('two_pointers_to_struct_members')
+  .desc(`Test aliasing of two struct member pointers passed to a function.`)
+  .params(u =>
+    u
+      .combine('address_space', kWritableAddressSpaces)
+      .combine('a_use', keysOf(kUses))
+      .combine('b_use', keysOf(kUses))
+      .combine('member', ['a', 'b'])
+      .combine('aliased', [true, false])
+      .beginSubcases()
+  )
+  .fn(t => {
+    t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+
+    const code = `
+struct S {
+  a : i32,
+  b : i32,
+}
+
+${maybeDeclareModuleScopeVar('x', t.params.address_space, 'S')}
+${maybeDeclareModuleScopeVar('y', t.params.address_space, 'S')}
+
+fn callee(pa : ${ptr(t.params.address_space, 'i32')},
+          pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
+  ${kUses[t.params.a_use].gen(`*pa`)}
+  ${kUses[t.params.b_use].gen(`*pb`)}
+  return 0;
+}
+
+fn caller() {
+  ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'S')}
+  ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'S')}
+  callee(&x.${t.params.member}, ${t.params.aliased ? `&x.a` : `&y.a`});
+}
+`;
+    t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code);
+  });
+
+g.test('two_pointers_to_struct_members_indirect')
+  .desc(
+    `Test aliasing of two structure pointers passed to a function, which accesses members of those
+structures and then passes the member pointers to another function.`
+  )
+  .params(u =>
+    u
+      .combine('address_space', kWritableAddressSpaces)
+      .combine('a_use', keysOf(kUses))
+      .combine('b_use', keysOf(kUses))
+      .combine('member', ['a', 'b'])
+      .combine('aliased', [true, false])
+      .beginSubcases()
+  )
+  .fn(t => {
+    t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+
+    const code = `
+struct S {
+  a : i32,
+  b : i32,
+}
+
+${maybeDeclareModuleScopeVar('x', t.params.address_space, 'S')}
+${maybeDeclareModuleScopeVar('y', t.params.address_space, 'S')}
+
+fn callee(pa : ${ptr(t.params.address_space, 'i32')},
+          pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
+  ${kUses[t.params.a_use].gen(`*pa`)}
+  ${kUses[t.params.b_use].gen(`*pb`)}
+  return 0;
+}
+
+fn access(pa : ${ptr(t.params.address_space, 'S')},
+          pb : ${ptr(t.params.address_space, 'S')}) -> i32 {
+  return callee(&(*pa).${t.params.member}, &(*pb).a);
+}
+
+fn caller() {
+  ${maybeDeclareFunctionScopeVar('x', t.params.address_space, 'S')}
+  ${maybeDeclareFunctionScopeVar('y', t.params.address_space, 'S')}
+  access(&x, ${t.params.aliased ? `&x` : `&y`});
+}
+`;
+    t.expectCompileResult(shouldPass(t.params.aliased, t.params.a_use, t.params.b_use), code);
+  });
+
 g.test('one_pointer_one_module_scope')
   .desc(`Test aliasing of a pointer with a direct access to a module-scope variable.`)
   .params(u =>
     u
+      .combine('address_space', ['private', 'storage', 'workgroup'] as const)
       .combine('a_use', keysOf(kUses))
       .combine('b_use', keysOf(kUses))
       .combine('aliased', [true, false])
       .beginSubcases()
   )
   .fn(t => {
+    if (requiresUnrestrictedPointerParameters(t.params.address_space)) {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+
     const code = `
-var<private> x : i32;
-var<private> y : i32;
+${declareModuleScopeVar('x', t.params.address_space, 'i32')}
+${declareModuleScopeVar('y', t.params.address_space, 'i32')}
 
-fn callee(pb : ptr<private, i32>) -> i32 {
+fn callee(pb : ${ptr(t.params.address_space, 'i32')}) -> i32 {
   ${kUses[t.params.a_use].gen(`x`)}
   ${kUses[t.params.b_use].gen(`*pb`)}
   return 0;
@@ -98,29 +332,34 @@ g.test('subcalls')
   .desc(`Test aliasing of two pointers passed to a function, and then passed to other functions.`)
   .params(u =>
     u
+      .combine('address_space', ['private', 'storage', 'workgroup'] as const)
       .combine('a_use', ['no_access', 'assign', 'binary_lhs'] as UseName[])
       .combine('b_use', ['no_access', 'assign', 'binary_lhs'] as UseName[])
       .combine('aliased', [true, false])
       .beginSubcases()
   )
   .fn(t => {
+    if (requiresUnrestrictedPointerParameters(t.params.address_space)) {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+    const ptr_i32 = ptr(t.params.address_space, 'i32');
     const code = `
-var<private> x : i32;
-var<private> y : i32;
+${declareModuleScopeVar('x', t.params.address_space, 'i32')}
+${declareModuleScopeVar('y', t.params.address_space, 'i32')}
 
-fn subcall_no_access(p : ptr<private, i32>) {
+fn subcall_no_access(p : ${ptr_i32}) {
   let pp = &*p;
 }
 
-fn subcall_binary_lhs(p : ptr<private, i32>) -> i32 {
+fn subcall_binary_lhs(p : ${ptr_i32}) -> i32 {
   return *p + 1;
 }
 
-fn subcall_assign(p : ptr<private, i32>) {
+fn subcall_assign(p : ${ptr_i32}) {
   *p = 42;
 }
 
-fn callee(pa : ptr<private, i32>, pb : ptr<private, i32>) -> i32 {
+fn callee(pa : ${ptr_i32}, pb : ${ptr_i32}) -> i32 {
   let new_pa = &*pa;
   let new_pb = &*pb;
   subcall_${t.params.a_use}(new_pa);
@@ -139,20 +378,25 @@ g.test('member_accessors')
   .desc(`Test aliasing of two pointers passed to a function and used with member accessors.`)
   .params(u =>
     u
+      .combine('address_space', ['private', 'storage', 'workgroup'] as const)
       .combine('a_use', ['no_access', 'assign', 'binary_lhs'] as UseName[])
       .combine('b_use', ['no_access', 'assign', 'binary_lhs'] as UseName[])
       .combine('aliased', [true, false])
       .beginSubcases()
   )
   .fn(t => {
+    if (requiresUnrestrictedPointerParameters(t.params.address_space)) {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+
+    const ptr_S = ptr(t.params.address_space, 'S');
     const code = `
 struct S { a : i32 }
 
-var<private> x : S;
-var<private> y : S;
+${declareModuleScopeVar('x', t.params.address_space, 'S')}
+${declareModuleScopeVar('y', t.params.address_space, 'S')}
 
-fn callee(pa : ptr<private, S>,
-          pb : ptr<private, S>) -> i32 {
+fn callee(pa : ${ptr_S}, pb : ${ptr_S}) -> i32 {
   ${kUses[t.params.a_use].gen(`(*pa).a`)}
   ${kUses[t.params.b_use].gen(`(*pb).a`)}
   return 0;
@@ -167,12 +411,18 @@ fn caller() {
 
 g.test('same_pointer_read_and_write')
   .desc(`Test that we can read from and write to the same pointer.`)
-  .params(u => u.beginSubcases())
+  .params(u =>
+    u.combine('address_space', ['private', 'storage', 'workgroup'] as const).beginSubcases()
+  )
   .fn(t => {
+    if (requiresUnrestrictedPointerParameters(t.params.address_space)) {
+      t.skipIfLanguageFeatureNotSupported('unrestricted_pointer_parameters');
+    }
+
     const code = `
-var<private> v : i32;
+${declareModuleScopeVar('v', t.params.address_space, 'i32')}
 
-fn callee(p : ptr<private, i32>) {
+fn callee(p : ${ptr(t.params.address_space, 'i32')}) {
   *p = *p + 1;
 }
 
@@ -185,10 +435,12 @@ fn caller() {
 
 g.test('aliasing_inside_function')
   .desc(`Test that we can alias pointers inside a function.`)
-  .params(u => u.beginSubcases())
+  .params(u =>
+    u.combine('address_space', ['private', 'storage', 'workgroup'] as const).beginSubcases()
+  )
   .fn(t => {
     const code = `
-var<private> v : i32;
+${declareModuleScopeVar('v', t.params.address_space, 'i32')}
 
 fn foo() {
   var v : i32;
diff --git a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts
index 41249e445d1c..aaba1f95aa54 100644
--- a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts
+++ b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts
@@ -367,7 +367,14 @@ function generatePointerCheck(check: string): string {
   }
 }
 
-const kPointerCases = {
+interface PointerCase {
+  code: string;
+  check: 'address' | 'contents';
+  uniform: boolean | 'never';
+  needs_deref_sugar?: boolean;
+}
+
+const kPointerCases: Record<string, PointerCase> = {
   address_uniform_literal: {
     code: `let ptr = &wg_array[0];`,
     check: `address`,
@@ -585,6 +592,168 @@ const kPointerCases = {
     check: `contents`,
     uniform: false,
   },
+  contents_lhs_ref_pointer_deref1: {
+    code: `*&func_scalar = uniform_value;
+    let test_val = func_scalar;`,
+    check: `contents`,
+    uniform: true,
+  },
+  contents_lhs_ref_pointer_deref1a: {
+    code: `*&func_scalar = nonuniform_value;
+    let test_val = func_scalar;`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref2: {
+    code: `*&(func_array[nonuniform_value]) = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref2a: {
+    code: `(func_array[nonuniform_value]) = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref3: {
+    code: `*&(func_array[needs_uniform(uniform_value)]) = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: true,
+  },
+  contents_lhs_ref_pointer_deref3a: {
+    code: `*&(func_array[needs_uniform(nonuniform_value)]) = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: 'never',
+  },
+  contents_lhs_ref_pointer_deref4: {
+    code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[uniform_value]) = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: true,
+  },
+  contents_lhs_ref_pointer_deref4a: {
+    code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[uniform_value]) = nonuniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref4b: {
+    code: `*&((*&(func_struct.x[uniform_value])).x[uniform_value].x[nonuniform_value]) = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref4c: {
+    code: `*&((*&(func_struct.x[uniform_value])).x[nonuniform_value]).x[uniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref4d: {
+    code: `*&((*&(func_struct.x[nonuniform_value])).x[uniform_value].x)[uniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+  },
+  contents_lhs_ref_pointer_deref4e: {
+    code: `*&((*&(func_struct.x[uniform_value])).x[needs_uniform(nonuniform_value)].x[uniform_value]) = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: 'never',
+  },
+
+  // The following cases require the 'pointer_composite_access' language feature.
+  contents_lhs_pointer_deref2: {
+    code: `(&func_array)[uniform_value] = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: true,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref2a: {
+    code: `(&func_array)[nonuniform_value] = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref3: {
+    code: `(&func_array)[needs_uniform(uniform_value)] = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: true,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref3a: {
+    code: `(&func_array)[needs_uniform(nonuniform_value)] = uniform_value;
+    let test_val = func_array[0];`,
+    check: `contents`,
+    uniform: 'never',
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4: {
+    code: `(&((&(func_struct.x[uniform_value])).x[uniform_value]).x)[uniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: true,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4a: {
+    code: `(&((&(func_struct.x[uniform_value])).x[uniform_value]).x)[uniform_value] = nonuniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4b: {
+    code: `(&((&(func_struct.x[uniform_value])).x)[uniform_value]).x[nonuniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4c: {
+    code: `(&((&(func_struct.x[uniform_value])).x[nonuniform_value]).x)[uniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4d: {
+    code: `(&((&(func_struct.x[nonuniform_value])).x[uniform_value]).x)[uniform_value] = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_lhs_pointer_deref4e: {
+    code: `(&((&(func_struct.x[uniform_value])).x)[needs_uniform(nonuniform_value)].x[uniform_value]) = uniform_value;
+    let test_val = func_struct.x[0].x[0].x[0];`,
+    check: `contents`,
+    uniform: 'never',
+    needs_deref_sugar: true,
+  },
+  contents_rhs_pointer_deref1: {
+    code: `let test_val = (&func_array)[uniform_value];`,
+    check: `contents`,
+    uniform: true,
+    needs_deref_sugar: true,
+  },
+  contents_rhs_pointer_deref1a: {
+    code: `let test_val = (&func_array)[nonuniform_value];`,
+    check: `contents`,
+    uniform: false,
+    needs_deref_sugar: true,
+  },
+  contents_rhs_pointer_deref2: {
+    code: `let test_val = (&func_array)[needs_uniform(nonuniform_value)];`,
+    check: `contents`,
+    uniform: `never`,
+    needs_deref_sugar: true,
+  },
 };
 
 g.test('pointers')
@@ -612,6 +781,13 @@ var<storage> uniform_value : u32;
 @group(0) @binding(1)
 var<storage, read_write> nonuniform_value : u32;
 
+fn needs_uniform(val : u32) -> u32{
+  if val == 0 {
+    workgroupBarrier();
+  }
+  return val;
+}
+
 @compute @workgroup_size(16, 1, 1)
 fn main(@builtin(local_invocation_id) lid : vec3<u32>,
         @builtin(global_invocation_id) gid : vec3<u32>) {
@@ -627,11 +803,16 @@ fn main(@builtin(local_invocation_id) lid : vec3<u32>,
       `
 ${generatePointerCheck(testcase.check)}
 }`;
-    if (!testcase.uniform) {
+
+    if (testcase.needs_deref_sugar === true) {
+      t.skipIfLanguageFeatureNotSupported('pointer_composite_access');
+    }
+    // Explicitly check false to distinguish from never.
+    if (testcase.uniform === false) {
       const without_check = code + `}\n`;
       t.expectCompileResult(true, without_check);
     }
-    t.expectCompileResult(testcase.uniform, with_check);
+    t.expectCompileResult(testcase.uniform === true, with_check);
   });
 
 function expectedUniformity(uniform: string, init: string): boolean {
diff --git a/src/webgpu/web_platform/canvas/readbackFromWebGPUCanvas.spec.ts b/src/webgpu/web_platform/canvas/readbackFromWebGPUCanvas.spec.ts
index 7fd7142f00e4..50b095d899e0 100644
--- a/src/webgpu/web_platform/canvas/readbackFromWebGPUCanvas.spec.ts
+++ b/src/webgpu/web_platform/canvas/readbackFromWebGPUCanvas.spec.ts
@@ -14,7 +14,12 @@ TODO: implement all canvas types, see TODO on kCanvasTypes.
 `;
 
 import { makeTestGroup } from '../../../common/framework/test_group.js';
-import { assert, raceWithRejectOnTimeout, unreachable } from '../../../common/util/util.js';
+import {
+  ErrorWithExtra,
+  assert,
+  raceWithRejectOnTimeout,
+  unreachable,
+} from '../../../common/util/util.js';
 import {
   kCanvasAlphaModes,
   kCanvasColorSpaces,
@@ -28,6 +33,8 @@ import {
   createCanvas,
   createOnscreenCanvas,
 } from '../../util/create_elements.js';
+import { TexelView } from '../../util/texture/texel_view.js';
+import { findFailedPixels } from '../../util/texture/texture_ok.js';
 
 export const g = makeTestGroup(GPUTest);
 
@@ -180,9 +187,40 @@ function readPixelsFrom2DCanvasAndCompare(
   ctx: CanvasRenderingContext2D | OffscreenCanvasRenderingContext2D,
   expect: Uint8ClampedArray
 ) {
-  const actual = ctx.getImageData(0, 0, ctx.canvas.width, ctx.canvas.height).data;
+  const { width, height } = ctx.canvas;
+  const actual = ctx.getImageData(0, 0, width, height).data;
+
+  const subrectOrigin = [0, 0, 0];
+  const subrectSize = [width, height, 1];
+
+  const areaDesc = {
+    bytesPerRow: width * 4,
+    rowsPerImage: height,
+    subrectOrigin,
+    subrectSize,
+  };
 
-  t.expectOK(checkElementsEqual(actual, expect));
+  const format = 'rgba8unorm';
+  const actTexelView = TexelView.fromTextureDataByReference(format, actual, areaDesc);
+  const expTexelView = TexelView.fromTextureDataByReference(format, expect, areaDesc);
+
+  const failedPixelsMessage = findFailedPixels(
+    format,
+    { x: 0, y: 0, z: 0 },
+    { width, height, depthOrArrayLayers: 1 },
+    { actTexelView, expTexelView },
+    { maxFractionalDiff: 0 }
+  );
+
+  if (failedPixelsMessage !== undefined) {
+    const msg = 'Canvas had unexpected contents:\n' + failedPixelsMessage;
+    t.expectOK(
+      new ErrorWithExtra(msg, () => ({
+        expTexelView,
+        actTexelView,
+      }))
+    );
+  }
 }
 
 g.test('onscreenCanvas,snapshot')
diff --git a/src/webgpu/web_platform/util.ts b/src/webgpu/web_platform/util.ts
index 358272a37927..20fdca59c638 100644
--- a/src/webgpu/web_platform/util.ts
+++ b/src/webgpu/web_platform/util.ts
@@ -295,6 +295,70 @@ export const kVideoInfo = makeTable({
         bottomRightColor: 'green',
       },
     },
+    'four-colors-h264-bt601-hflip.mp4': {
+      mimeType: 'video/mp4; codecs=vp09.00.10.08',
+      colorSpace: 'bt601',
+      coded: {
+        topLeftColor: 'yellow',
+        topRightColor: 'red',
+        bottomLeftColor: 'blue',
+        bottomRightColor: 'green',
+      },
+      display: {
+        topLeftColor: 'red',
+        topRightColor: 'yellow',
+        bottomLeftColor: 'green',
+        bottomRightColor: 'blue',
+      },
+    },
+    'four-colors-h264-bt601-vflip.mp4': {
+      mimeType: 'video/mp4; codecs=vp09.00.10.08',
+      colorSpace: 'bt601',
+      coded: {
+        topLeftColor: 'yellow',
+        topRightColor: 'red',
+        bottomLeftColor: 'blue',
+        bottomRightColor: 'green',
+      },
+      display: {
+        topLeftColor: 'blue',
+        topRightColor: 'green',
+        bottomLeftColor: 'yellow',
+        bottomRightColor: 'red',
+      },
+    },
+    'four-colors-vp9-bt601-hflip.mp4': {
+      mimeType: 'video/mp4; codecs=vp09.00.10.08',
+      colorSpace: 'bt601',
+      coded: {
+        topLeftColor: 'yellow',
+        topRightColor: 'red',
+        bottomLeftColor: 'blue',
+        bottomRightColor: 'green',
+      },
+      display: {
+        topLeftColor: 'red',
+        topRightColor: 'yellow',
+        bottomLeftColor: 'green',
+        bottomRightColor: 'blue',
+      },
+    },
+    'four-colors-vp9-bt601-vflip.mp4': {
+      mimeType: 'video/mp4; codecs=vp09.00.10.08',
+      colorSpace: 'bt601',
+      coded: {
+        topLeftColor: 'yellow',
+        topRightColor: 'red',
+        bottomLeftColor: 'blue',
+        bottomRightColor: 'green',
+      },
+      display: {
+        topLeftColor: 'blue',
+        topRightColor: 'green',
+        bottomLeftColor: 'yellow',
+        bottomRightColor: 'red',
+      },
+    },
   },
 } as const);