diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts
index fddb65e1420..88f51dfdfe0 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts
@@ -32,6 +32,7 @@ import {
   kDepthStencilFormats,
   kEncodableTextureFormats,
 } from '../../../../../format_info.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   appendComponentTypeForFormatToTextureType,
@@ -83,6 +84,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isFillable(t.format))
       .combine('minFilter', ['nearest', 'linear'] as const)
@@ -98,7 +100,8 @@ Parameters:
     skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format);
   })
   .fn(async t => {
-    const { format, C, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, C, samplePoints, stage, addressModeU, addressModeV, minFilter, offset } =
+      t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -137,14 +140,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -169,6 +181,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isFillable(t.format))
       .combine('minFilter', ['nearest', 'linear'] as const)
@@ -182,7 +195,7 @@ Parameters:
     skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format);
   })
   .fn(async t => {
-    const { format, C, samplePoints, addressMode, minFilter } = t.params;
+    const { format, C, stage, samplePoints, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube';
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -225,14 +238,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = appendComponentTypeForFormatToTextureType('texture_cube', format);
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -266,6 +288,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isFillable(t.format))
       .combine('minFilter', ['nearest', 'linear'] as const)
@@ -282,7 +305,8 @@ Parameters:
     skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, C, A, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, C, A, addressModeU, addressModeV, minFilter, offset } =
+      t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -326,14 +350,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d_array', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -361,6 +394,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isFillable(t.format))
       .combine('minFilter', ['nearest', 'linear'] as const)
@@ -376,7 +410,7 @@ Parameters:
     skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format);
   })
   .fn(async t => {
-    const { format, C, A, samplePoints, addressMode, minFilter } = t.params;
+    const { format, C, A, stage, samplePoints, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -421,14 +455,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = appendComponentTypeForFormatToTextureType('texture_cube_array', format);
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -454,6 +497,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -467,7 +511,7 @@ Parameters:
       .combine('offset', [false, true] as const)
   )
   .fn(async t => {
-    const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -503,14 +547,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -529,6 +582,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -540,7 +594,7 @@ Parameters:
       .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const)
   )
   .fn(async t => {
-    const { format, samplePoints, addressMode, minFilter } = t.params;
+    const { format, stage, samplePoints, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube';
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -580,14 +634,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -616,6 +679,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -634,7 +698,8 @@ Parameters:
     skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, A, addressModeU, addressModeV, minFilter, offset } =
+      t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -675,14 +740,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d_array';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -704,6 +778,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -719,7 +794,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported('cube-array');
   })
   .fn(async t => {
-    const { format, A, samplePoints, addressMode, minFilter } = t.params;
+    const { format, A, stage, samplePoints, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -761,14 +836,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube_array';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts
index cad0e85c662..742badde94a 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts
@@ -23,6 +23,7 @@ import {
   isEncodableTextureFormat,
   kDepthStencilFormats,
 } from '../../../../../format_info.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   checkCallResults,
@@ -67,6 +68,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -85,8 +87,17 @@ Parameters:
     t.skipIfTextureFormatNotSupported(t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, A, addressModeU, addressModeV, minFilter, compare, offset } =
-      t.params;
+    const {
+      format,
+      stage,
+      samplePoints,
+      A,
+      addressModeU,
+      addressModeV,
+      minFilter,
+      compare,
+      offset,
+    } = t.params;
 
     const viewDimension = '2d-array';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension });
@@ -130,14 +141,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d_array';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -160,6 +180,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -176,7 +197,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported('cube-array');
   })
   .fn(async t => {
-    const { format, A, samplePoints, addressMode, minFilter, compare } = t.params;
+    const { format, A, stage, samplePoints, addressMode, minFilter, compare } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -222,14 +243,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube_array';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -256,6 +286,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -270,7 +301,7 @@ Parameters:
       .combine('offset', [false, true] as const)
   )
   .fn(async t => {
-    const { format, C, samplePoints, addressMode, compare, minFilter, offset } = t.params;
+    const { format, C, stage, samplePoints, addressMode, compare, minFilter, offset } = t.params;
 
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
     const descriptor: GPUTextureDescriptor = {
@@ -309,14 +340,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -336,6 +376,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -348,7 +389,7 @@ Parameters:
       .combine('compare', kCompareFunctions)
   )
   .fn(async t => {
-    const { format, samplePoints, addressMode, minFilter, compare } = t.params;
+    const { format, stage, samplePoints, addressMode, minFilter, compare } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -391,14 +432,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts
index 2f42fffeb76..9bb4f74b035 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts
@@ -38,6 +38,7 @@ import {
 import { GPUTest } from '../../../../../gpu_test.js';
 import { maxMipLevelCount, virtualMipSize } from '../../../../../util/texture/base.js';
 import { TexelFormats } from '../../../../types.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   TextureCall,
@@ -90,6 +91,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => textureDimensionAndFormatCompatible('1d', t.format))
       // 1d textures can't have a height !== 1
@@ -105,7 +107,7 @@ Parameters:
     t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format);
   })
   .fn(async t => {
-    const { format, C, L, samplePoints } = t.params;
+    const { format, stage, C, L, samplePoints } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -137,14 +139,23 @@ Parameters:
     const textureType = appendComponentTypeForFormatToTextureType('texture_1d', texture.format);
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -166,6 +177,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => !isCompressedFloatTextureFormat(t.format))
       .beginSubcases()
@@ -179,7 +191,7 @@ Parameters:
     t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, C, L } = t.params;
+    const { format, stage, samplePoints, C, L } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -209,14 +221,23 @@ Parameters:
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d', texture.format);
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -237,6 +258,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => textureDimensionAndFormatCompatible('3d', t.format))
       .beginSubcases()
@@ -250,7 +272,7 @@ Parameters:
     t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, C, L } = t.params;
+    const { format, stage, samplePoints, C, L } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' });
@@ -282,14 +304,23 @@ Parameters:
     const textureType = appendComponentTypeForFormatToTextureType('texture_3d', texture.format);
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -312,6 +343,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('texture_type', [
         'texture_multisampled_2d',
         'texture_depth_multisampled_2d',
@@ -336,7 +368,7 @@ Parameters:
     t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format);
   })
   .fn(async t => {
-    const { texture_type, format, samplePoints, C, S } = t.params;
+    const { texture_type, format, stage, samplePoints, C, S } = t.params;
 
     const sampleCount = 4;
     const descriptor: GPUTextureDescriptor = {
@@ -368,14 +400,23 @@ Parameters:
     const textureType = appendComponentTypeForFormatToTextureType(texture_type, texture.format);
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -396,6 +437,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -410,7 +452,7 @@ Parameters:
     t.skipIfTextureLoadNotSupportedForTextureType('texture_depth_2d');
   })
   .fn(async t => {
-    const { format, samplePoints, C, L } = t.params;
+    const { format, stage, samplePoints, C, L } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -440,14 +482,23 @@ Parameters:
     const textureType = 'texture_depth_2d';
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -465,8 +516,10 @@ Parameters:
  * coords: The 0-based texel coordinate.
 `
   )
-  .paramsSubcasesOnly(u =>
+  .params(u =>
     u
+      .combine('stage', kShaderStages)
+      .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
       .combine('C', ['i32', 'u32'] as const)
       .combine('L', ['i32', 'u32'] as const)
@@ -475,7 +528,7 @@ Parameters:
     t.skipIf(typeof VideoFrame === 'undefined', 'VideoFrames are not supported')
   )
   .fn(async t => {
-    const { samplePoints, C, L } = t.params;
+    const { stage, samplePoints, C, L } = t.params;
 
     const size = [8, 8, 1];
 
@@ -505,14 +558,23 @@ Parameters:
     const textureType = 'texture_external';
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
     videoFrame.close();
@@ -536,6 +598,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       // MAINTENANCE_TODO: Update createTextureFromTexelViews to support stencil8 and remove this filter.
       .filter(t => t.format !== 'stencil8' && !isCompressedFloatTextureFormat(t.format))
@@ -559,7 +622,7 @@ Parameters:
     t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format);
   })
   .fn(async t => {
-    const { texture_type, format, samplePoints, C, A, L } = t.params;
+    const { texture_type, format, stage, samplePoints, C, A, L } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' });
@@ -595,14 +658,23 @@ Parameters:
     const textureType = appendComponentTypeForFormatToTextureType(texture_type, texture.format);
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -622,6 +694,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const)
       .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
@@ -636,7 +709,7 @@ Parameters:
     }
   })
   .fn(async t => {
-    const { format, samplePoints, C } = t.params;
+    const { format, stage, samplePoints, C } = t.params;
 
     // We want at least 3 blocks or something wide enough for 3 mip levels.
     const [width] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -663,14 +736,23 @@ Parameters:
     const textureType = `texture_storage_1d<${format}, read>`;
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -690,6 +772,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const)
       .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
@@ -704,7 +787,7 @@ Parameters:
     }
   })
   .fn(async t => {
-    const { format, samplePoints, C } = t.params;
+    const { format, stage, samplePoints, C } = t.params;
 
     // We want at least 3 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 3, format });
@@ -729,14 +812,23 @@ Parameters:
     const textureType = `texture_storage_2d<${format}, read>`;
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -758,6 +850,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const)
       .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
@@ -773,7 +866,7 @@ Parameters:
     }
   })
   .fn(async t => {
-    const { format, samplePoints, C, A } = t.params;
+    const { format, stage, samplePoints, C, A } = t.params;
 
     // We want at least 3 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' });
@@ -803,14 +896,23 @@ Parameters:
       dimension: '2d-array',
     };
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -830,6 +932,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const)
       .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
@@ -844,7 +947,7 @@ Parameters:
     }
   })
   .fn(async t => {
-    const { format, samplePoints, C } = t.params;
+    const { format, stage, samplePoints, C } = t.params;
 
     // We want at least 3 blocks or something wide enough for 3 mip levels.
     const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' });
@@ -870,14 +973,23 @@ Parameters:
     const textureType = `texture_storage_3d<${format}, read>`;
     const viewDescriptor = {};
     const sampler = undefined;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts
index b670c44035f..e3bf2c01b13 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts
@@ -5,6 +5,7 @@ Execution tests for textureSampleBaseClampToEdge
 import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
 import { GPUTest } from '../../../../../gpu_test.js';
 import { TexelView } from '../../../../../util/texture/texel_view.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   checkCallResults,
@@ -54,6 +55,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('textureType', ['texture_2d<f32>', 'texture_external'] as const)
       .beginSubcases()
       .combine('samplePoints', kSamplePointMethods)
@@ -68,7 +70,7 @@ Parameters:
     )
   )
   .fn(async t => {
-    const { textureType, samplePoints, addressModeU, addressModeV, minFilter } = t.params;
+    const { textureType, stage, samplePoints, addressModeU, addressModeV, minFilter } = t.params;
 
     const descriptor: GPUTextureDescriptor = {
       format: 'rgba8unorm',
@@ -105,14 +107,23 @@ Parameters:
         };
       });
       const viewDescriptor = {};
-      const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+      const results = await doTextureCalls(
+        t,
+        texture,
+        viewDescriptor,
+        textureType,
+        sampler,
+        calls,
+        stage
+      );
       const res = await checkCallResults(
         t,
         { texels, descriptor, viewDescriptor },
         textureType,
         sampler,
         calls,
-        results
+        results,
+        stage
       );
       t.expectOK(res);
     } finally {
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts
index 75077b5afa7..1062a76a4c7 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts
@@ -18,6 +18,7 @@ import {
   isEncodableTextureFormat,
   kDepthStencilFormats,
 } from '../../../../../format_info.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   checkCallResults,
@@ -59,6 +60,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -73,7 +75,7 @@ Parameters:
       .combine('offset', [false, true] as const)
   )
   .fn(async t => {
-    const { format, samplePoints, addressModeU, addressModeV, minFilter, compare, offset } =
+    const { format, stage, samplePoints, addressModeU, addressModeV, minFilter, compare, offset } =
       t.params;
 
     const size = chooseTextureSize({ minSize: 16, minBlocks: 4, format });
@@ -117,14 +119,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -144,6 +155,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -156,7 +168,7 @@ Parameters:
       .combine('compare', kCompareFunctions)
   )
   .fn(async t => {
-    const { format, samplePoints, addressMode, minFilter, compare } = t.params;
+    const { format, stage, samplePoints, addressMode, minFilter, compare } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube';
     const size = chooseTextureSize({ minSize: 16, minBlocks: 2, format, viewDimension });
@@ -203,14 +215,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -240,6 +261,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -258,8 +280,17 @@ Parameters:
     t.skipIfTextureFormatNotSupported(t.params.format);
   })
   .fn(async t => {
-    const { format, samplePoints, A, addressModeU, addressModeV, minFilter, compare, offset } =
-      t.params;
+    const {
+      format,
+      stage,
+      samplePoints,
+      A,
+      addressModeU,
+      addressModeV,
+      minFilter,
+      compare,
+      offset,
+    } = t.params;
 
     const viewDimension = '2d-array';
     const size = chooseTextureSize({ minSize: 16, minBlocks: 4, format, viewDimension });
@@ -306,14 +337,23 @@ Parameters:
     });
     const textureType = 'texture_depth_2d_array';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -336,6 +376,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -352,7 +393,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported('cube-array');
   })
   .fn(async t => {
-    const { format, A, samplePoints, addressMode, minFilter, compare } = t.params;
+    const { format, A, stage, samplePoints, addressMode, minFilter, compare } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension });
@@ -400,14 +441,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = 'texture_depth_cube_array';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts
index bbff69d3296..774fe0f16cc 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts
@@ -7,6 +7,7 @@ Samples a texture using explicit gradients.
 
 import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
 import { kCompressedTextureFormats, kEncodableTextureFormats } from '../../../../../format_info.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   appendComponentTypeForFormatToTextureType,
@@ -57,6 +58,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -70,7 +72,7 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -108,14 +110,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -144,6 +155,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .combine('viewDimension', ['3d', 'cube'] as const)
@@ -165,6 +177,7 @@ Parameters:
     const {
       format,
       viewDimension,
+      stage,
       samplePoints,
       addressModeU,
       addressModeV,
@@ -240,14 +253,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = getTextureTypeForTextureViewDimension(viewDimension)!;
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -278,6 +300,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -292,7 +315,8 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, A, addressModeU, addressModeV, minFilter, offset } =
+      t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -335,14 +359,23 @@ Parameters:
     });
     const textureType = 'texture_2d_array<f32>';
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -372,6 +405,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -385,7 +419,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported('cube-array');
   })
   .fn(async t => {
-    const { format, samplePoints, A, addressMode, minFilter } = t.params;
+    const { format, stage, samplePoints, A, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({
@@ -433,14 +467,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = getTextureTypeForTextureViewDimension(viewDimension);
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts
index 5c6e99eb966..14f6b936a42 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts
@@ -1,9 +1,6 @@
 export const description = `
 Samples a texture.
 
-Must only be used in a fragment shader stage.
-Must only be invoked in uniform control flow.
-
 - TODO: Test un-encodable formats.
 `;
 
@@ -16,6 +13,7 @@ import {
   kDepthStencilFormats,
   kEncodableTextureFormats,
 } from '../../../../../format_info.js';
+import { kShaderStages } from '../../../../validation/decl/util.js';
 
 import {
   appendComponentTypeForFormatToTextureType,
@@ -69,6 +67,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -82,7 +81,7 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -120,14 +119,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -161,6 +169,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -175,7 +184,8 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params;
+    const { format, stage, samplePoints, A, addressModeU, addressModeV, minFilter, offset } =
+      t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -218,14 +228,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_2d_array', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -257,6 +276,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .combine('viewDimension', ['3d', 'cube'] as const)
@@ -273,7 +293,7 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, viewDimension, samplePoints, addressMode, minFilter, offset } = t.params;
+    const { format, viewDimension, stage, samplePoints, addressMode, minFilter, offset } = t.params;
 
     const [width, height] = chooseTextureSize({ minSize: 32, minBlocks: 2, format, viewDimension });
     const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension);
@@ -327,14 +347,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = getTextureTypeForTextureViewDimension(viewDimension);
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -369,6 +398,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kTestableColorFormats)
       .filter(t => isPotentiallyFilterableAndFillable(t.format))
       .beginSubcases()
@@ -382,7 +412,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported('cube-array');
   })
   .fn(async t => {
-    const { format, samplePoints, A, addressMode, minFilter } = t.params;
+    const { format, stage, samplePoints, A, addressMode, minFilter } = t.params;
 
     const viewDimension: GPUTextureViewDimension = 'cube-array';
     const size = chooseTextureSize({
@@ -429,14 +459,23 @@ Parameters:
       dimension: viewDimension,
     };
     const textureType = getTextureTypeForTextureViewDimension(viewDimension);
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -469,6 +508,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -485,7 +525,7 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, addressMode, minFilter, L, offset } = t.params;
+    const { format, stage, samplePoints, addressMode, minFilter, L, offset } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -523,14 +563,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_depth_2d', format);
     const viewDescriptor = {};
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -565,6 +614,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -582,7 +632,7 @@ Parameters:
     skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format)
   )
   .fn(async t => {
-    const { format, samplePoints, addressMode, minFilter, A, L, offset } = t.params;
+    const { format, stage, samplePoints, addressMode, minFilter, A, L, offset } = t.params;
 
     // We want at least 4 blocks or something wide enough for 3 mip levels.
     const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format });
@@ -624,14 +674,23 @@ Parameters:
     });
     const textureType = appendComponentTypeForFormatToTextureType('texture_depth_2d_array', format);
     const viewDescriptor: GPUTextureViewDescriptor = { dimension: '2d-array' };
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
     const res = await checkCallResults(
       t,
       { texels, descriptor, viewDescriptor },
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
@@ -665,6 +724,7 @@ Parameters:
   )
   .params(u =>
     u
+      .combine('stage', kShaderStages)
       .combine('format', kDepthStencilFormats)
       // filter out stencil only formats
       .filter(t => isDepthTextureFormat(t.format))
@@ -686,7 +746,7 @@ Parameters:
     t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension);
   })
   .fn(async t => {
-    const { format, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params;
+    const { format, stage, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params;
 
     const size = chooseTextureSize({
       minSize: 32,
@@ -734,7 +794,15 @@ Parameters:
     };
     const textureType =
       viewDimension === 'cube' ? 'texture_depth_cube' : 'texture_depth_cube_array';
-    const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls);
+    const results = await doTextureCalls(
+      t,
+      texture,
+      viewDescriptor,
+      textureType,
+      sampler,
+      calls,
+      stage
+    );
 
     const res = await checkCallResults(
       t,
@@ -742,7 +810,8 @@ Parameters:
       textureType,
       sampler,
       calls,
-      results
+      results,
+      stage
     );
     t.expectOK(res);
   });
diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
index 2f64739c86c..b706a7add37 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts
@@ -40,6 +40,7 @@ import {
 import { PerPixelAtLevel, TexelView } from '../../../../../util/texture/texel_view.js';
 import { createTextureFromTexelViews } from '../../../../../util/texture.js';
 import { reifyExtent3D } from '../../../../../util/unions.js';
+import { ShaderStage } from '../../../../validation/decl/util.js';
 
 export const kSampleTypeInfo = {
   f32: {
@@ -164,22 +165,66 @@ type MipWeights = {
 };
 type MipWeightType = keyof MipWeights;
 
+function makeGraph(width: number, height: number) {
+  const data = new Uint8Array(width * height);
+
+  return {
+    plot(norm: number, x: number, c: number) {
+      const y = clamp(Math.round(norm * height), { min: 0, max: height - 1 });
+      const offset = (height - y - 1) * width + x;
+      data[offset] = c;
+    },
+    plotValues(values: Iterable<number>, c: number) {
+      let i = 0;
+      for (const v of values) {
+        this.plot(v, i, c);
+        ++i;
+      }
+    },
+    toString(conversion = ['.', 'e', 'A']) {
+      const lines = [];
+      for (let y = 0; y < height; ++y) {
+        const offset = y * width;
+        lines.push([...data.subarray(offset, offset + width)].map(v => conversion[v]).join(''));
+      }
+      return lines.join('\n');
+    },
+  };
+}
+
+function* linear0to1OverN(n: number) {
+  for (let i = 0; i <= n; ++i) {
+    yield i / n;
+  }
+}
+
+function graphWeights(height: number, weights: number[]) {
+  const graph = makeGraph(weights.length, height);
+  graph.plotValues(linear0to1OverN(weights.length - 1), 1);
+  graph.plotValues(weights, 2);
+  return graph.toString();
+}
+
 /**
  * Validates the weights go from 0 to 1 in increasing order.
  */
-function validateWeights(weights: number[]) {
-  const showWeights = () => weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n');
+function validateWeights(stage: string, weights: number[]) {
+  const showWeights = () => `
+${weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n')}
+
+e = expected
+A = actual
+${graphWeights(32, weights)}
+`;
 
   // Validate the weights
-  assert(weights[0] === 0, `weight 0 expected 0 but was ${weights[0]}\n${showWeights()}`);
   assert(
-    weights[kMipGradientSteps] === 1,
-    `top weight expected 1 but was ${weights[kMipGradientSteps]}\n${showWeights()}`
+    weights[0] === 0,
+    `stage: ${stage}, weight 0 expected 0 but was ${weights[0]}\n${showWeights()}`
   );
-
   assert(
-    new Set(weights).size >= ((weights.length * 0.66) | 0),
-    `expected more unique weights\n${showWeights()}`
+    weights[kMipGradientSteps] === 1,
+    `stage: ${stage}, top weight expected 1 but was ${weights[kMipGradientSteps]}\n${showWeights()}`
   );
 
   // Note: for 16 steps, these are the AMD weights
@@ -207,12 +252,68 @@ function validateWeights(weights: number[]) {
   //
   // notice step 1 is 0 and step 15 is 1.
   // so we only check the 1 through 14.
+  //
+  // Note: these 2 changes are effectively here to catch Intel Mac
+  // issues and require implementations to work around them.
+  //
+  // Ideally the weights should form a straight line
+  //
+  // +----------------+
+  // |              **|
+  // |            **  |
+  // |          **    |
+  // |        **      |
+  // |      **        |
+  // |    **          |
+  // |  **            |
+  // |**              |
+  // +----------------+
+  //
+  // AMD Mac goes like this: Not great but we allow it
+  //
+  // +----------------+
+  // |             ***|
+  // |           **   |
+  // |          *     |
+  // |        **      |
+  // |      **        |
+  // |     *          |
+  // |   **           |
+  // |***             |
+  // +----------------+
+  //
+  // Intel Mac goes like this: Unacceptable
+  //
+  // +----------------+
+  // |         *******|
+  // |         *      |
+  // |        *       |
+  // |        *       |
+  // |       *        |
+  // |       *        |
+  // |      *         |
+  // |*******         |
+  // +----------------+
+  //
+  const dx = 1 / kMipGradientSteps;
   for (let i = 0; i < kMipGradientSteps; ++i) {
+    const dy = weights[i + 1] - weights[i];
+    // dy / dx because dy might be 0
+    const slope = dy / dx;
+    assert(
+      slope >= 0,
+      `stage: ${stage}, weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}`
+    );
     assert(
-      weights[i] <= weights[i + 1],
-      `weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}`
+      slope <= 2,
+      `stage: ${stage}, slope from weight[${i}] to weight[${i + 1}] is > 2.\n${showWeights()}`
     );
   }
+
+  assert(
+    new Set(weights).size >= ((weights.length * 0.66) | 0),
+    `stage: ${stage}, expected more unique weights\n${showWeights()}`
+  );
 }
 
 /**
@@ -328,43 +429,78 @@ async function queryMipGradientValuesForDevice(t: GPUTest) {
       struct VSOutput {
         @builtin(position) pos: vec4f,
         @location(0) @interpolate(flat, either) ndx: u32,
+        @location(1) @interpolate(flat, either) result: vec4f,
       };
 
-      @vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput {
+      fn getMixLevels(wNdx: u32) -> vec4f {
+        let mipLevel = f32(wNdx) / ${kMipGradientSteps};
+        let size = textureDimensions(tex);
+        let g = mix(1.0, 2.0, mipLevel) / f32(size.x);
+        let ddx = vec2f(g, 0);
+        return vec4f(
+          textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r,
+          textureSampleGrad(tex, smp, vec2f(0.5), ddx, vec2f(0)).r,
+          0,
+          0);
+      }
+
+      fn recordMixLevels(wNdx: u32, r: vec4f) {
+        let ndx = wNdx * ${kNumWeightTypes};
+        for (var i: u32 = 0; i < ${kNumWeightTypes}; i++) {
+          result[ndx + i] = r[i];
+        }
+      }
+
+      fn getPosition(vNdx: u32) -> vec4f {
         let pos = array(
           vec2f(-1,  3),
           vec2f( 3, -1),
           vec2f(-1, -1),
         );
         let p = pos[vNdx];
-        return VSOutput(vec4f(p, 0, 1), iNdx);
+        return vec4f(p, 0, 1);
       }
 
-      @fragment fn fs(v: VSOutput) -> @location(0) vec4f {
-        let mipLevel = f32(v.ndx) / ${kMipGradientSteps};
-        let size = textureDimensions(tex);
-        let g = mix(1.0, 2.0, mipLevel) / f32(size.x);
-        let ddx = vec2f(g, 0);
+      @vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput {
+        return VSOutput(getPosition(vNdx), iNdx, vec4f(0));
+      }
+
+      @fragment fn fsRecord(v: VSOutput) -> @location(0) vec4f {
+        recordMixLevels(v.ndx, getMixLevels(v.ndx));
+        return vec4f(0);
+      }
+
+      @compute @workgroup_size(1) fn csRecord(@builtin(global_invocation_id) id: vec3u) {
+        recordMixLevels(id.x, getMixLevels(id.x));
+      }
+
+      @vertex fn vsRecord(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput {
+        return VSOutput(getPosition(vNdx), iNdx, getMixLevels(iNdx));
 
-        let ndx = v.ndx * ${kNumWeightTypes};
-        result[ndx + 0] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r;
-        result[ndx + 1] = textureSampleGrad(tex, smp, vec2f(0.5), ddx, vec2f(0)).r;
+      }
 
+      @fragment fn fsSaveVs(v: VSOutput) -> @location(0) vec4f {
+        recordMixLevels(v.ndx, v.result);
         return vec4f(0);
       }
     `,
   });
 
-  const pipeline = device.createRenderPipeline({
+  const vertexPipeline = device.createRenderPipeline({
     layout: 'auto',
-    vertex: { module },
-    fragment: { module, targets: [{ format: 'rgba8unorm' }] },
+    vertex: { module, entryPoint: 'vsRecord' },
+    fragment: { module, entryPoint: 'fsSaveVs', targets: [{ format: 'rgba8unorm' }] },
   });
 
-  const target = t.createTextureTracked({
-    size: [1, 1],
-    format: 'rgba8unorm',
-    usage: GPUTextureUsage.RENDER_ATTACHMENT,
+  const fragmentPipeline = device.createRenderPipeline({
+    layout: 'auto',
+    vertex: { module, entryPoint: 'vs' },
+    fragment: { module, entryPoint: 'fsRecord', targets: [{ format: 'rgba8unorm' }] },
+  });
+
+  const computePipeline = device.createComputePipeline({
+    layout: 'auto',
+    compute: { module },
   });
 
   const texture = t.createTextureTracked({
@@ -387,59 +523,131 @@ async function queryMipGradientValuesForDevice(t: GPUTest) {
     mipmapFilter: 'linear',
   });
 
+  const target = t.createTextureTracked({
+    size: [1, 1],
+    format: 'rgba8unorm',
+    usage: GPUTextureUsage.RENDER_ATTACHMENT,
+  });
+
   const storageBuffer = t.createBufferTracked({
     size: 4 * (kMipGradientSteps + 1) * kNumWeightTypes,
     usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
   });
 
-  const resultBuffer = t.createBufferTracked({
-    size: storageBuffer.size,
-    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
-  });
+  type PassFn = (
+    encoder: GPUCommandEncoder,
+    bindGroup: GPUBindGroup,
+    resultBuffer: GPUBuffer
+  ) => void;
+
+  const getMixWeightForStage = (
+    encoder: GPUCommandEncoder,
+    pipeline: GPUComputePipeline | GPURenderPipeline,
+    passFn: PassFn
+  ) => {
+    const resultBuffer = t.createBufferTracked({
+      size: storageBuffer.size,
+      usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
+    });
 
-  const bindGroup = device.createBindGroup({
-    layout: pipeline.getBindGroupLayout(0),
-    entries: [
-      { binding: 0, resource: texture.createView() },
-      { binding: 1, resource: sampler },
-      { binding: 2, resource: { buffer: storageBuffer } },
-    ],
-  });
+    const bindGroup = device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: texture.createView() },
+        { binding: 1, resource: sampler },
+        { binding: 2, resource: { buffer: storageBuffer } },
+      ],
+    });
+
+    passFn(encoder, bindGroup, resultBuffer);
+    encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size);
+    return resultBuffer;
+  };
 
   const encoder = device.createCommandEncoder();
-  const pass = encoder.beginRenderPass({
-    colorAttachments: [
-      {
-        view: target.createView(),
-        loadOp: 'clear',
-        storeOp: 'store',
-      },
-    ],
-  });
-  pass.setPipeline(pipeline);
-  pass.setBindGroup(0, bindGroup);
-  pass.draw(3, kMipGradientSteps + 1);
-  pass.end();
-  encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size);
+  const stageBuffers: Record<ShaderStage, GPUBuffer> = {
+    compute: getMixWeightForStage(
+      encoder,
+      computePipeline,
+      (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => {
+        const pass = encoder.beginComputePass();
+        pass.setPipeline(computePipeline);
+        pass.setBindGroup(0, bindGroup);
+        pass.dispatchWorkgroups(kMipGradientSteps + 1);
+        pass.end();
+      }
+    ),
+    fragment: getMixWeightForStage(
+      encoder,
+      fragmentPipeline,
+      (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => {
+        const pass = encoder.beginRenderPass({
+          colorAttachments: [
+            {
+              view: target.createView(),
+              loadOp: 'clear',
+              storeOp: 'store',
+            },
+          ],
+        });
+        pass.setPipeline(fragmentPipeline);
+        pass.setBindGroup(0, bindGroup);
+        pass.draw(3, kMipGradientSteps + 1);
+        pass.end();
+      }
+    ),
+    vertex: getMixWeightForStage(
+      encoder,
+      vertexPipeline,
+      (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => {
+        const pass = encoder.beginRenderPass({
+          colorAttachments: [
+            {
+              view: target.createView(),
+              loadOp: 'clear',
+              storeOp: 'store',
+            },
+          ],
+        });
+        pass.setPipeline(vertexPipeline);
+        pass.setBindGroup(0, bindGroup);
+        pass.draw(kMipGradientSteps + 1);
+        pass.end();
+      }
+    ),
+  };
   device.queue.submit([encoder.finish()]);
 
-  await resultBuffer.mapAsync(GPUMapMode.READ);
-  const result = Array.from(new Float32Array(resultBuffer.getMappedRange()));
-  resultBuffer.unmap();
+  await Promise.all(Object.values(stageBuffers).map(b => b.mapAsync(GPUMapMode.READ)));
 
-  const [sampleLevelWeights, gradWeights] = unzip(result, kNumWeightTypes);
+  const mixWeightsByStage = Object.fromEntries(
+    Object.entries(stageBuffers).map(([stage, resultBuffer]) => {
+      const result = Array.from(new Float32Array(resultBuffer.getMappedRange()));
+      resultBuffer.unmap();
+      resultBuffer.destroy();
 
-  validateWeights(sampleLevelWeights);
-  validateWeights(gradWeights);
+      const [sampleLevelWeights, gradWeights] = unzip(result, kNumWeightTypes);
 
-  s_deviceToMipGradientValues.set(device, {
-    sampleLevelWeights,
-    softwareMixToGPUMixGradWeights: generateSoftwareMixToGPUMixGradWeights(gradWeights, texture),
-  });
+      validateWeights(stage, sampleLevelWeights);
+      validateWeights(stage, gradWeights);
+
+      return [
+        stage,
+        {
+          sampleLevelWeights,
+          softwareMixToGPUMixGradWeights: generateSoftwareMixToGPUMixGradWeights(
+            gradWeights,
+            texture
+          ),
+        },
+      ];
+    })
+  ) as Record<ShaderStage, MipWeights>;
+
+  s_deviceToMipGradientValues.set(device, mixWeightsByStage);
 
   texture.destroy();
   storageBuffer.destroy();
-  resultBuffer.destroy();
 }
 
 // Given an array of ascending values and a value v, finds
@@ -511,10 +719,15 @@ function generateSoftwareMixToGPUMixGradWeights(gpuWeights: number[], texture: G
   return softwareMixToGPUMixMap;
 }
 
-function mapSoftwareMipLevelToGPUMipLevel(t: GPUTest, mipLevel: number) {
+function mapSoftwareMipLevelToGPUMipLevel(t: GPUTest, stage: ShaderStage, mipLevel: number) {
   const baseLevel = Math.floor(mipLevel);
   const softwareMix = mipLevel - baseLevel;
-  const gpuMix = getMixWeightByTypeForMipLevel(t, 'softwareMixToGPUMixGradWeights', softwareMix);
+  const gpuMix = getMixWeightByTypeForMipLevel(
+    t,
+    stage,
+    'softwareMixToGPUMixGradWeights',
+    softwareMix
+  );
   return baseLevel + gpuMix;
 }
 
@@ -549,7 +762,7 @@ const euclideanModulo = (n: number, m: number) => ((n % m) + m) % m;
  */
 const kMipGradientSteps = 64;
 const s_deviceToMipGradientValuesPromise = new WeakMap<GPUDevice, Promise<void>>();
-const s_deviceToMipGradientValues = new WeakMap<GPUDevice, MipWeights>();
+const s_deviceToMipGradientValues = new WeakMap<GPUDevice, Record<ShaderStage, MipWeights>>();
 async function initMipGradientValuesForDevice(t: GPUTest) {
   const { device } = t;
   let weightsP = s_deviceToMipGradientValuesPromise.get(device);
@@ -562,6 +775,7 @@ async function initMipGradientValuesForDevice(t: GPUTest) {
 
 function getMixWeightByTypeForMipLevel(
   t: GPUTest,
+  stage: ShaderStage,
   weightType: MipWeightType | 'identity',
   mipLevel: number
 ) {
@@ -569,7 +783,7 @@ function getMixWeightByTypeForMipLevel(
     return euclideanModulo(mipLevel, 1);
   }
   // linear interpolate between weights
-  const weights = s_deviceToMipGradientValues.get(t.device)![weightType];
+  const weights = s_deviceToMipGradientValues.get(t.device)![stage][weightType];
   assert(
     !!weights,
     'you must use WGSLTextureSampleTest or call initializeDeviceMipWeights before calling this function'
@@ -584,6 +798,7 @@ function getMixWeightByTypeForMipLevel(
 
 function getWeightForMipLevel(
   t: GPUTest,
+  stage: ShaderStage,
   weightType: MipWeightType | 'identity',
   mipLevelCount: number,
   mipLevel: number
@@ -591,7 +806,7 @@ function getWeightForMipLevel(
   if (mipLevel < 0 || mipLevel >= mipLevelCount) {
     return 1;
   }
-  return getMixWeightByTypeForMipLevel(t, weightType, mipLevel);
+  return getMixWeightByTypeForMipLevel(t, stage, weightType, mipLevel);
 }
 
 /**
@@ -1435,6 +1650,7 @@ function softwareTextureReadMipLevel<T extends Dimensionality>(
  */
 function softwareTextureReadLevel<T extends Dimensionality>(
   t: GPUTest,
+  stage: ShaderStage,
   call: TextureCall<T>,
   texture: Texture,
   sampler: GPUSamplerDescriptor | undefined,
@@ -1456,7 +1672,7 @@ function softwareTextureReadLevel<T extends Dimensionality>(
       const t0 = softwareTextureReadMipLevel<T>(call, texture, sampler, baseMipLevel);
       const t1 = softwareTextureReadMipLevel<T>(call, texture, sampler, nextMipLevel);
       const weightType = call.builtin === 'textureSampleLevel' ? 'sampleLevelWeights' : 'identity';
-      const mix = getWeightForMipLevel(t, weightType, mipLevelCount, clampedMipLevel);
+      const mix = getWeightForMipLevel(t, stage, weightType, mipLevelCount, clampedMipLevel);
       assert(mix >= 0 && mix <= 1);
       const values = [
         { v: t0, weight: 1 - mix },
@@ -1518,6 +1734,7 @@ function computeMipLevelFromGradientsForCall<T extends Dimensionality>(
  */
 function softwareTextureReadGrad<T extends Dimensionality>(
   t: GPUTest,
+  stage: ShaderStage,
   call: TextureCall<T>,
   texture: Texture,
   sampler?: GPUSamplerDescriptor
@@ -1525,10 +1742,10 @@ function softwareTextureReadGrad<T extends Dimensionality>(
   const bias = call.bias === undefined ? 0 : clamp(call.bias, { min: -16.0, max: 15.99 });
   if (call.ddx) {
     const mipLevel = computeMipLevelFromGradientsForCall(call, texture.descriptor.size);
-    const weightMipLevel = mapSoftwareMipLevelToGPUMipLevel(t, mipLevel + bias);
-    return softwareTextureReadLevel(t, call, texture, sampler, weightMipLevel);
+    const weightMipLevel = mapSoftwareMipLevelToGPUMipLevel(t, stage, mipLevel + bias);
+    return softwareTextureReadLevel(t, stage, call, texture, sampler, weightMipLevel);
   } else {
-    return softwareTextureReadLevel(t, call, texture, sampler, (call.mipLevel ?? 0) + bias);
+    return softwareTextureReadLevel(t, stage, call, texture, sampler, (call.mipLevel ?? 0) + bias);
   }
 }
 
@@ -1599,6 +1816,7 @@ function derivativeForCall<T extends Dimensionality>(
 
 function softwareTextureRead<T extends Dimensionality>(
   t: GPUTest,
+  stage: ShaderStage,
   call: TextureCall<T>,
   texture: Texture,
   sampler?: GPUSamplerDescriptor
@@ -1612,7 +1830,7 @@ function softwareTextureRead<T extends Dimensionality>(
     };
     call = newCall;
   }
-  return softwareTextureReadGrad(t, call, texture, sampler);
+  return softwareTextureReadGrad(t, stage, call, texture, sampler);
 }
 
 export type TextureTestOptions<T extends Dimensionality> = {
@@ -1841,7 +2059,8 @@ export async function checkCallResults<T extends Dimensionality>(
   textureType: string,
   sampler: GPUSamplerDescriptor | undefined,
   calls: TextureCall<T>[],
-  results: Awaited<ReturnType<typeof doTextureCalls<T>>>
+  results: Awaited<ReturnType<typeof doTextureCalls<T>>>,
+  stage: ShaderStage = 'fragment' // MAINTENANCE_TODO: remove default
 ) {
   const errs: string[] = [];
   const format = texture.texels[0].format;
@@ -1856,7 +2075,7 @@ export async function checkCallResults<T extends Dimensionality>(
   for (let callIdx = 0; callIdx < calls.length; callIdx++) {
     const call = calls[callIdx];
     const gotRGBA = results.results[callIdx];
-    const expectRGBA = softwareTextureRead(t, call, texture, sampler);
+    const expectRGBA = softwareTextureRead(t, stage, call, texture, sampler);
 
     // The spec says depth and stencil have implementation defined values for G, B, and A
     // so if this is `textureGather` and component > 0 then there's nothing to check.
@@ -1975,6 +2194,7 @@ export async function checkCallResults<T extends Dimensionality>(
               return Promise.resolve(
                 softwareTextureRead(
                   t,
+                  stage,
                   call,
                   {
                     texels,
@@ -3863,7 +4083,10 @@ function describeTextureCall<T extends Dimensionality>(call: TextureCall<T>): st
   return `${call.builtin}(${args.join(', ')})`;
 }
 
-const s_deviceToPipelines = new WeakMap<GPUDevice, Map<string, GPURenderPipeline>>();
+const s_deviceToPipelines = new WeakMap<
+  GPUDevice,
+  Map<string, GPURenderPipeline | GPUComputePipeline>
+>();
 
 /**
  * Given a list of "calls", each one of which has a texture coordinate,
@@ -3937,7 +4160,8 @@ export async function doTextureCalls<T extends Dimensionality>(
   viewDescriptor: GPUTextureViewDescriptor,
   textureType: string,
   sampler: GPUSamplerDescriptor | undefined,
-  calls: TextureCall<T>[]
+  calls: TextureCall<T>[],
+  stage: ShaderStage = 'fragment' // MAINTENANCE_TODO: remove default
 ) {
   const {
     format,
@@ -4021,6 +4245,51 @@ export async function doTextureCalls<T extends Dimensionality>(
       ? 'vec3f(v.pos.xy - 0.5, 0) / vec3f(textureDimensions(T))'
       : '(v.pos.xy - 0.5) / vec2f(textureDimensions(T))'
   };`;
+  const derivativeType =
+    isCubeViewDimension(viewDescriptor) || dimension === '3d'
+      ? 'vec3f'
+      : dimension === '1d'
+      ? 'f32'
+      : 'vec2f';
+
+  const stageWGSL =
+    stage === 'vertex'
+      ? `
+// --------------------------- vertex stage shaders --------------------------------
+@vertex fn vsVertex(
+    @builtin(vertex_index) vertex_index : u32,
+    @builtin(instance_index) instance_index : u32) -> VOut {
+  let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
+  return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, getResult(instance_index, ${derivativeType}(0)));
+}
+
+@fragment fn fsVertex(v: VOut) -> @location(0) vec4f {
+  results[v.ndx] = v.result;
+  return vec4f(0);
+}
+`
+      : stage === 'fragment'
+      ? `
+// --------------------------- fragment stage shaders --------------------------------
+@vertex fn vsFragment(
+    @builtin(vertex_index) vertex_index : u32,
+    @builtin(instance_index) instance_index : u32) -> VOut {
+  let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
+  return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0));
+}
+
+@fragment fn fsFragment(v: VOut) -> @location(0) vec4f {
+  ${derivativeBaseWGSL}
+  results[v.ndx] = getResult(v.ndx, derivativeBase);
+  return vec4f(0);
+}
+`
+      : `
+// --------------------------- compute stage shaders --------------------------------
+@compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) {
+  results[id.x] = getResult(id.x, ${derivativeType}(0));
+}
+`;
 
   const code = `
 ${structs}
@@ -4032,33 +4301,25 @@ ${dataFields}
 struct VOut {
   @builtin(position) pos: vec4f,
   @location(0) @interpolate(flat, either) ndx: u32,
+  @location(1) @interpolate(flat, either) result: ${returnType},
 };
 
-@vertex
-fn vs_main(
-    @builtin(vertex_index) vertex_index : u32,
-    @builtin(instance_index) instance_index : u32) -> VOut {
-  let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1));
-  return VOut(vec4f(positions[vertex_index], 0, 1), instance_index);
-}
-
 @group(0) @binding(0) var          T    : ${textureType};
 ${sampler ? `@group(0) @binding(1) var          S    : ${samplerType}` : ''};
 @group(0) @binding(2) var<storage> data : Data;
-@group(0) @binding(3) var<storage, read_write> results: array<${returnType}>;
+@group(1) @binding(0) var<storage, read_write> results: array<${returnType}>;
 
-@fragment
-fn fs_main(v: VOut) -> @location(0) vec4f {
-  ${derivativeBaseWGSL}
-  let idx = v.ndx;
+fn getResult(idx: u32, derivativeBase: ${derivativeType}) -> ${returnType} {
   var result : ${resultType};
 ${body}
-  results[idx] = ${returnType}(result);
-  return vec4f(0);
+  return ${returnType}(result);
 }
+
+${stageWGSL}
 `;
 
-  const pipelines = s_deviceToPipelines.get(t.device) ?? new Map<string, GPURenderPipeline>();
+  const pipelines =
+    s_deviceToPipelines.get(t.device) ?? new Map<string, GPURenderPipeline | GPUComputePipeline>();
   s_deviceToPipelines.set(t.device, pipelines);
 
   // unfilterable-float textures can only be used with manually created bindGroupLayouts
@@ -4088,21 +4349,21 @@ ${body}
     sampleType = 'unfilterable-float';
   }
 
+  const visibility =
+    stage === 'compute'
+      ? GPUShaderStage.COMPUTE
+      : stage === 'fragment'
+      ? GPUShaderStage.FRAGMENT
+      : GPUShaderStage.VERTEX;
+
   const entries: GPUBindGroupLayoutEntry[] = [
     {
       binding: 2,
-      visibility: GPUShaderStage.FRAGMENT,
+      visibility,
       buffer: {
         type: 'read-only-storage',
       },
     },
-    {
-      binding: 3,
-      visibility: GPUShaderStage.FRAGMENT,
-      buffer: {
-        type: 'storage',
-      },
-    },
   ];
 
   const viewDimension = effectiveViewDimensionForDimension(
@@ -4114,7 +4375,7 @@ ${body}
   if (textureType.includes('storage')) {
     entries.push({
       binding: 0,
-      visibility: GPUShaderStage.FRAGMENT,
+      visibility,
       storageTexture: {
         access: 'read-only',
         viewDimension,
@@ -4124,13 +4385,13 @@ ${body}
   } else if (gpuTexture instanceof GPUExternalTexture) {
     entries.push({
       binding: 0,
-      visibility: GPUShaderStage.FRAGMENT,
+      visibility,
       externalTexture: {},
     });
   } else {
     entries.push({
       binding: 0,
-      visibility: GPUShaderStage.FRAGMENT,
+      visibility,
       texture: {
         sampleType,
         viewDimension,
@@ -4142,31 +4403,64 @@ ${body}
   if (sampler) {
     entries.push({
       binding: 1,
-      visibility: GPUShaderStage.FRAGMENT,
+      visibility,
       sampler: {
         type: isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering',
       },
     });
   }
 
-  const id = `${resultType}:${JSON.stringify(entries)}:${code}`;
+  const id = `${resultType}:${stage}:${JSON.stringify(entries)}:${code}`;
   let pipeline = pipelines.get(id);
   if (!pipeline) {
-    const shaderModule = t.device.createShaderModule({ code });
-    const bindGroupLayout = t.device.createBindGroupLayout({ entries });
-    const layout = t.device.createPipelineLayout({
-      bindGroupLayouts: [bindGroupLayout],
+    const module = t.device.createShaderModule({ code });
+    const bindGroupLayout0 = t.device.createBindGroupLayout({ entries });
+    const bindGroupLayout1 = t.device.createBindGroupLayout({
+      entries: [
+        {
+          binding: 0,
+          visibility: GPUShaderStage.FRAGMENT | GPUShaderStage.COMPUTE,
+          buffer: {
+            type: 'storage',
+          },
+        },
+      ],
     });
 
-    pipeline = t.device.createRenderPipeline({
-      layout,
-      vertex: { module: shaderModule },
-      fragment: {
-        module: shaderModule,
-        targets: [{ format: renderTarget.format }],
-      },
+    const layout = t.device.createPipelineLayout({
+      bindGroupLayouts: [bindGroupLayout0, bindGroupLayout1],
     });
 
+    switch (stage) {
+      case 'compute':
+        pipeline = t.device.createComputePipeline({
+          layout,
+          compute: { module },
+        });
+        break;
+      case 'fragment':
+        pipeline = t.device.createRenderPipeline({
+          layout,
+          vertex: { module, entryPoint: 'vsFragment' },
+          fragment: {
+            module,
+            entryPoint: 'fsFragment',
+            targets: [{ format: renderTarget.format }],
+          },
+        });
+        break;
+      case 'vertex':
+        pipeline = t.device.createRenderPipeline({
+          layout,
+          vertex: { module, entryPoint: 'vsVertex' },
+          fragment: {
+            module,
+            entryPoint: 'fsVertex',
+            targets: [{ format: renderTarget.format }],
+          },
+        });
+        break;
+    }
     pipelines.set(id, pipeline);
   }
 
@@ -4178,7 +4472,7 @@ ${body}
       usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
     });
 
-    const bindGroup = t.device.createBindGroup({
+    const bindGroup0 = t.device.createBindGroup({
       layout: pipeline!.getBindGroupLayout(0),
       entries: [
         {
@@ -4190,10 +4484,14 @@ ${body}
         },
         ...(sampler ? [{ binding: 1, resource: gpuSampler! }] : []),
         { binding: 2, resource: { buffer: dataBuffer } },
-        { binding: 3, resource: { buffer: storageBuffer } },
       ],
     });
 
+    const bindGroup1 = t.device.createBindGroup({
+      layout: pipeline!.getBindGroupLayout(1),
+      entries: [{ binding: 0, resource: { buffer: storageBuffer } }],
+    });
+
     const resultBuffer = t.createBufferTracked({
       size: storageBuffer.size,
       usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
@@ -4201,20 +4499,30 @@ ${body}
 
     const encoder = t.device.createCommandEncoder();
 
-    const renderPass = encoder.beginRenderPass({
-      colorAttachments: [
-        {
-          view: renderTarget.createView(),
-          loadOp: 'clear',
-          storeOp: 'store',
-        },
-      ],
-    });
+    if (stage === 'compute') {
+      const pass = encoder.beginComputePass();
+      pass.setPipeline(pipeline! as GPUComputePipeline);
+      pass.setBindGroup(0, bindGroup0);
+      pass.setBindGroup(1, bindGroup1);
+      pass.dispatchWorkgroups(calls.length);
+      pass.end();
+    } else {
+      const pass = encoder.beginRenderPass({
+        colorAttachments: [
+          {
+            view: renderTarget.createView(),
+            loadOp: 'clear',
+            storeOp: 'store',
+          },
+        ],
+      });
 
-    renderPass.setPipeline(pipeline!);
-    renderPass.setBindGroup(0, bindGroup);
-    renderPass.draw(3, calls.length);
-    renderPass.end();
+      pass.setPipeline(pipeline! as GPURenderPipeline);
+      pass.setBindGroup(0, bindGroup0);
+      pass.setBindGroup(1, bindGroup1);
+      pass.draw(3, calls.length);
+      pass.end();
+    }
     encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size);
     t.device.queue.submit([encoder.finish()]);