diff --git a/src/webgpu/api/validation/layout_shader_compat.spec.ts b/src/webgpu/api/validation/layout_shader_compat.spec.ts
index 986fc4229655..2b5e609c55d6 100644
--- a/src/webgpu/api/validation/layout_shader_compat.spec.ts
+++ b/src/webgpu/api/validation/layout_shader_compat.spec.ts
@@ -1,14 +1,293 @@
 export const description = `
 TODO:
 - interface matching between pipeline layout and shader
-    - x= {compute, vertex, fragment, vertex+fragment}, visibilities
     - x= bind group index values, binding index values, multiple bindings
-    - x= types of bindings
-    - x= {equal, superset, subset}
+    - x= {superset, subset}
 `;
 
 import { makeTestGroup } from '../../../common/framework/test_group.js';
+import {
+  kShaderStageCombinations,
+  kShaderStages,
+  ValidBindableResource,
+} from '../../capability_info.js';
+import { GPUConst } from '../../constants.js';
 
 import { ValidationTest } from './validation_test.js';
 
-export const g = makeTestGroup(ValidationTest);
+type BindableResourceType = ValidBindableResource | 'readonlyStorageBuf';
+const kBindableResources = [
+  'uniformBuf',
+  'storageBuf',
+  'readonlyStorageBuf',
+  'filtSamp',
+  'nonFiltSamp',
+  'compareSamp',
+  'sampledTex',
+  'sampledTexMS',
+  'readonlyStorageTex',
+  'writeonlyStorageTex',
+  'readwriteStorageTex',
+] as const;
+
+const bindGroupLayoutEntryContents = {
+  compareSamp: {
+    sampler: {
+      type: 'comparison',
+    },
+  },
+  filtSamp: {
+    sampler: {
+      type: 'filtering',
+    },
+  },
+  nonFiltSamp: {
+    sampler: {
+      type: 'non-filtering',
+    },
+  },
+  sampledTex: {
+    texture: {
+      sampleType: 'unfilterable-float',
+    },
+  },
+  sampledTexMS: {
+    texture: {
+      sampleType: 'unfilterable-float',
+      multisampled: true,
+    },
+  },
+  storageBuf: {
+    buffer: {
+      type: 'storage',
+    },
+  },
+  readonlyStorageBuf: {
+    buffer: {
+      type: 'read-only-storage',
+    },
+  },
+  uniformBuf: {
+    buffer: {
+      type: 'uniform',
+    },
+  },
+  readonlyStorageTex: {
+    storageTexture: {
+      format: 'r32float',
+      access: 'read-only',
+    },
+  },
+  writeonlyStorageTex: {
+    storageTexture: {
+      format: 'r32float',
+      access: 'write-only',
+    },
+  },
+  readwriteStorageTex: {
+    storageTexture: {
+      format: 'r32float',
+      access: 'read-write',
+    },
+  },
+} as const;
+
+class F extends ValidationTest {
+  createPipelineLayout(
+    bindingInPipelineLayout: BindableResourceType,
+    visibility: number
+  ): GPUPipelineLayout {
+    return this.device.createPipelineLayout({
+      bindGroupLayouts: [
+        this.device.createBindGroupLayout({
+          entries: [
+            {
+              binding: 0,
+              visibility,
+              ...bindGroupLayoutEntryContents[bindingInPipelineLayout],
+            },
+          ],
+        }),
+      ],
+    });
+  }
+
+  GetBindableResourceShaderDeclaration(bindableResource: BindableResourceType): string {
+    switch (bindableResource) {
+      case 'compareSamp':
+        return 'var tmp : sampler_comparison';
+      case 'filtSamp':
+      case 'nonFiltSamp':
+        return 'var tmp : sampler';
+      case 'sampledTex':
+        return 'var tmp : texture_2d<f32>';
+      case 'sampledTexMS':
+        return 'var tmp : texture_multisampled_2d<f32>';
+      case 'storageBuf':
+        return 'var<storage, read_write> tmp : vec4u';
+      case 'readonlyStorageBuf':
+        return 'var<storage, read> tmp : vec4u';
+      case 'uniformBuf':
+        return 'var<uniform> tmp : vec4u;';
+      case 'readonlyStorageTex':
+        return 'var tmp : texture_storage_2d<r32float, read>';
+      case 'writeonlyStorageTex':
+        return 'var tmp : texture_storage_2d<r32float, write>';
+      case 'readwriteStorageTex':
+        return 'var tmp : texture_storage_2d<r32float, read_write>';
+    }
+  }
+}
+
+const BindingResourceCompatibleWithShaderStages = function (
+  bindingResource: BindableResourceType,
+  shaderStages: number
+): boolean {
+  if ((shaderStages & GPUConst.ShaderStage.VERTEX) > 0) {
+    switch (bindingResource) {
+      case 'writeonlyStorageTex':
+      case 'readwriteStorageTex':
+      case 'storageBuf':
+        return false;
+      default:
+        break;
+    }
+  }
+  return true;
+};
+
+export const g = makeTestGroup(F);
+
+g.test('pipeline_layout_shader_exact_match')
+  .desc(
+    `
+  Test that the binding type in the pipeline layout must match the related declaration in shader.
+  Note that read-write storage textures in the pipeline layout can match write-only storage textures
+  in the shader.
+  `
+  )
+  .params(u =>
+    u
+      .combine('bindingInPipelineLayout', kBindableResources)
+      .combine('bindingInShader', kBindableResources)
+      .beginSubcases()
+      .combine('pipelineLayoutVisibility', kShaderStageCombinations)
+      .combine('shaderStageWithBinding', kShaderStages)
+      .combine('isBindingStaticallyUsed', [true, false] as const)
+      .unless(
+        p =>
+          // We don't test using non-filtering sampler in shader because it has the same declaration
+          // as filtering sampler.
+          p.bindingInShader === 'nonFiltSamp' ||
+          !BindingResourceCompatibleWithShaderStages(
+            p.bindingInPipelineLayout,
+            p.pipelineLayoutVisibility
+          ) ||
+          !BindingResourceCompatibleWithShaderStages(p.bindingInShader, p.shaderStageWithBinding)
+      )
+  )
+  .fn(t => {
+    const {
+      bindingInPipelineLayout,
+      bindingInShader,
+      pipelineLayoutVisibility,
+      shaderStageWithBinding,
+      isBindingStaticallyUsed,
+    } = t.params;
+
+    const layout = t.createPipelineLayout(bindingInPipelineLayout, pipelineLayoutVisibility);
+    const bindResourceDeclaration = `@group(0) @binding(0) ${t.GetBindableResourceShaderDeclaration(
+      bindingInShader
+    )}`;
+    const staticallyUseBinding = isBindingStaticallyUsed ? '_ = tmp; ' : '';
+    const isAsync = false;
+
+    let success = true;
+    if (isBindingStaticallyUsed) {
+      success = bindingInPipelineLayout === bindingInShader;
+
+      // Filtering and non-filtering both have the same shader declaration.
+      success ||= bindingInPipelineLayout === 'nonFiltSamp' && bindingInShader === 'filtSamp';
+
+      // Promoting storage textures that are read-write in the layout can be readonly in the shader.
+      success ||=
+        bindingInPipelineLayout === 'readwriteStorageTex' &&
+        bindingInShader === 'writeonlyStorageTex';
+
+      // The shader using the resource must be included in the visibility in the layout.
+      success &&= (pipelineLayoutVisibility & shaderStageWithBinding) > 0;
+    }
+
+    switch (shaderStageWithBinding) {
+      case GPUConst.ShaderStage.COMPUTE: {
+        const computeShader = `
+        ${bindResourceDeclaration};
+        @compute @workgroup_size(1)
+        fn main() {
+          ${staticallyUseBinding}
+        }
+        `;
+        t.doCreateComputePipelineTest(isAsync, success, {
+          layout,
+          compute: {
+            module: t.device.createShaderModule({
+              code: computeShader,
+            }),
+          },
+        });
+        break;
+      }
+      case GPUConst.ShaderStage.VERTEX: {
+        const vertexShader = `
+        ${bindResourceDeclaration};
+        @vertex
+        fn main() -> @builtin(position) vec4f {
+          ${staticallyUseBinding}
+          return vec4f();
+        }
+        `;
+        t.doCreateRenderPipelineTest(isAsync, success, {
+          layout,
+          vertex: {
+            module: t.device.createShaderModule({
+              code: vertexShader,
+            }),
+          },
+        });
+        break;
+      }
+      case GPUConst.ShaderStage.FRAGMENT: {
+        const fragmentShader = `
+        ${bindResourceDeclaration};
+        @fragment
+        fn main() -> @location(0) vec4f {
+          ${staticallyUseBinding}
+          return vec4f();
+        }
+        `;
+        t.doCreateRenderPipelineTest(isAsync, success, {
+          layout,
+          vertex: {
+            module: t.device.createShaderModule({
+              code: `
+                @vertex
+                fn main() -> @builtin(position) vec4f {
+                  return vec4f();
+                }`,
+            }),
+          },
+          fragment: {
+            module: t.device.createShaderModule({
+              code: fragmentShader,
+            }),
+            targets: [
+              {
+                format: 'rgba8unorm',
+              },
+            ],
+          },
+        });
+        break;
+      }
+    }
+  });
diff --git a/src/webgpu/compat/api/validation/createBindGroupLayout.spec.ts b/src/webgpu/compat/api/validation/createBindGroupLayout.spec.ts
new file mode 100644
index 000000000000..f05af2860ec5
--- /dev/null
+++ b/src/webgpu/compat/api/validation/createBindGroupLayout.spec.ts
@@ -0,0 +1,34 @@
+export const description = `
+Tests that, in compat mode, you can not create a bind group layout with unsupported storage texture formats.
+`;
+
+import { makeTestGroup } from '../../../../common/framework/test_group.js';
+import { kCompatModeUnsupportedStorageTextureFormats } from '../../../format_info.js';
+import { CompatibilityTest } from '../../compatibility_test.js';
+
+export const g = makeTestGroup(CompatibilityTest);
+
+g.test('unsupportedStorageTextureFormats')
+  .desc(
+    `
+      Tests that, in compat mode, you can not create a bind group layout with unsupported storage texture formats.
+    `
+  )
+  .params(u => u.combine('format', kCompatModeUnsupportedStorageTextureFormats))
+  .fn(t => {
+    const { format } = t.params;
+
+    t.expectValidationError(() => {
+      t.device.createBindGroupLayout({
+        entries: [
+          {
+            binding: 0,
+            visibility: GPUShaderStage.COMPUTE,
+            storageTexture: {
+              format,
+            },
+          },
+        ],
+      });
+    }, true);
+  });
diff --git a/src/webgpu/compat/api/validation/render_pipeline/shader_module.spec.ts b/src/webgpu/compat/api/validation/render_pipeline/shader_module.spec.ts
index b89559b6fd10..fa6251afcd65 100644
--- a/src/webgpu/compat/api/validation/render_pipeline/shader_module.spec.ts
+++ b/src/webgpu/compat/api/validation/render_pipeline/shader_module.spec.ts
@@ -3,6 +3,7 @@ Tests limitations of createRenderPipeline related to shader modules in compat mo
 `;
 
 import { makeTestGroup } from '../../../../../common/framework/test_group.js';
+import { kCompatModeUnsupportedStorageTextureFormats } from '../../../../format_info.js';
 import { CompatibilityTest } from '../../../compatibility_test.js';
 
 export const g = makeTestGroup(CompatibilityTest);
@@ -152,3 +153,70 @@ Tests that you can not create a render pipeline with a shader module that uses i
       !isValid
     );
   });
+
+g.test('unsupportedStorageTextureFormats,computePipeline')
+  .desc(
+    `
+Tests that you can not create a compute pipeline with unsupported storage texture formats in compat mode.
+    `
+  )
+  .params(u =>
+    u //
+      .combine('format', kCompatModeUnsupportedStorageTextureFormats)
+      .combine('async', [false, true] as const)
+  )
+  .fn(t => {
+    const { format, async } = t.params;
+
+    const module = t.device.createShaderModule({
+      code: `
+        @group(0) @binding(0) var s: texture_storage_2d<${format}, read>;
+        @compute @workgroup_size(1) fn cs() {
+            _ = textureLoad(s, vec2u(0));
+        }
+      `,
+    });
+
+    const pipelineDescriptor: GPUComputePipelineDescriptor = {
+      layout: 'auto',
+      compute: {
+        module,
+        entryPoint: 'cs',
+      },
+    };
+    t.doCreateComputePipelineTest(async, false, pipelineDescriptor);
+  });
+
+g.test('unsupportedStorageTextureFormats,renderPipeline')
+  .desc(
+    `
+Tests that you can not create a render pipeline with unsupported storage texture formats in compat mode.
+    `
+  )
+  .params(u =>
+    u //
+      .combine('format', kCompatModeUnsupportedStorageTextureFormats)
+      .combine('async', [false, true] as const)
+  )
+  .fn(t => {
+    const { format, async } = t.params;
+
+    const module = t.device.createShaderModule({
+      code: `
+        @group(0) @binding(0) var s: texture_storage_2d<${format}, read>;
+        @vertex fn vs() -> @builtin(position) vec4f {
+            _ = textureLoad(s, vec2u(0));
+            return vec4f(0);
+        }
+      `,
+    });
+
+    const pipelineDescriptor: GPURenderPipelineDescriptor = {
+      layout: 'auto',
+      vertex: {
+        module,
+        entryPoint: 'vs',
+      },
+    };
+    t.doCreateRenderPipelineTest(async, false, pipelineDescriptor);
+  });
diff --git a/src/webgpu/compat/api/validation/texture/createTexture.spec.ts b/src/webgpu/compat/api/validation/texture/createTexture.spec.ts
index 4ce2d4556ae7..58dcd41ec764 100644
--- a/src/webgpu/compat/api/validation/texture/createTexture.spec.ts
+++ b/src/webgpu/compat/api/validation/texture/createTexture.spec.ts
@@ -6,7 +6,7 @@ Tests that textureBindingViewDimension must compatible with texture dimension
 
 import { makeTestGroup } from '../../../../../common/framework/test_group.js';
 import { kTextureDimensions, kTextureViewDimensions } from '../../../../capability_info.js';
-import { kColorTextureFormats, kTextureFormatInfo } from '../../../../format_info.js';
+import { kColorTextureFormats, kCompatModeUnsupportedStorageTextureFormats, kTextureFormatInfo } from '../../../../format_info.js';
 import { getTextureDimensionFromView } from '../../../../util/texture/base.js';
 import { CompatibilityTest } from '../../../compatibility_test.js';
 
@@ -153,3 +153,20 @@ g.test('format_reinterpretation')
       );
     }
   });
+
+g.test('unsupportedStorageTextureFormats')
+  .desc(`Tests that you can not create unsupported storage texture formats in compat mode.`)
+  .params(u => u.combine('format', kCompatModeUnsupportedStorageTextureFormats))
+  .fn(t => {
+    const { format } = t.params;
+    t.expectGPUError(
+      'validation',
+      () =>
+        t.device.createTexture({
+          size: [1, 1, 1],
+          format,
+          usage: GPUTextureUsage.STORAGE_BINDING,
+        }),
+      true
+    );
+  });
diff --git a/src/webgpu/format_info.ts b/src/webgpu/format_info.ts
index b8de86bb0c8b..64ef2cd8956d 100644
--- a/src/webgpu/format_info.ts
+++ b/src/webgpu/format_info.ts
@@ -1931,16 +1931,19 @@ export function isCompressedTextureFormat(format: GPUTextureFormat) {
   return format in kCompressedTextureFormatInfo;
 }
 
+export const kCompatModeUnsupportedStorageTextureFormats: readonly GPUTextureFormat[] = [
+  'rg32float',
+  'rg32sint',
+  'rg32uint',
+] as const;
+
 export function isTextureFormatUsableAsStorageFormat(
   format: GPUTextureFormat,
   isCompatibilityMode: boolean
 ) {
   if (isCompatibilityMode) {
-    switch (format) {
-      case 'rg32float':
-      case 'rg32sint':
-      case 'rg32uint':
-        return false;
+    if (kCompatModeUnsupportedStorageTextureFormats.indexOf(format) >= 0) {
+      return false;
     }
   }
   return !!kTextureFormatInfo[format].color?.storage;
diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json
index 213ca1d5b245..3c93ff996c11 100644
--- a/src/webgpu/listing_meta.json
+++ b/src/webgpu/listing_meta.json
@@ -602,6 +602,7 @@
   "webgpu:api,validation,image_copy,texture_related:texture,device_mismatch:*": { "subcaseMS": 5.417 },
   "webgpu:api,validation,image_copy,texture_related:usage:*": { "subcaseMS": 1.224 },
   "webgpu:api,validation,image_copy,texture_related:valid:*": { "subcaseMS": 3.678 },
+  "webgpu:api,validation,layout_shader_compat:pipeline_layout_shader_exact_match:*": { "subcaseMS": 2.00 },
   "webgpu:api,validation,query_set,create:count:*": { "subcaseMS": 0.967 },
   "webgpu:api,validation,query_set,destroy:invalid_queryset:*": { "subcaseMS": 0.801 },
   "webgpu:api,validation,query_set,destroy:twice:*": { "subcaseMS": 0.700 },
@@ -835,6 +836,7 @@
   "webgpu:api,validation,texture,rg11b10ufloat_renderable:create_render_pipeline:*": { "subcaseMS": 2.400 },
   "webgpu:api,validation,texture,rg11b10ufloat_renderable:create_texture:*": { "subcaseMS": 12.700 },
   "webgpu:compat,api,validation,createBindGroup:viewDimension_matches_textureBindingViewDimension:*": { "subcaseMS": 6.523 },
+  "webgpu:compat,api,validation,createBindGroupLayout:unsupportedStorageTextureFormats:*": { "subcaseMS": 0.601 },
   "webgpu:compat,api,validation,encoding,cmds,copyTextureToBuffer:compressed:*": { "subcaseMS": 202.929 },
   "webgpu:compat,api,validation,encoding,cmds,copyTextureToTexture:compressed:*": { "subcaseMS": 0.600 },
   "webgpu:compat,api,validation,encoding,programmable,pipeline_bind_group_compat:twoDifferentTextureViews,compute_pass,unused:*": { "subcaseMS": 1.501 },
@@ -844,10 +846,13 @@
   "webgpu:compat,api,validation,render_pipeline,fragment_state:colorState:*": { "subcaseMS": 32.604 },
   "webgpu:compat,api,validation,render_pipeline,shader_module:interpolate:*": { "subcaseMS": 1.502 },
   "webgpu:compat,api,validation,render_pipeline,shader_module:sample_mask:*": { "subcaseMS": 14.801 },
+  "webgpu:compat,api,validation,render_pipeline,shader_module:unsupportedStorageTextureFormats,computePipeline:*": { "subcaseMS": 0.601 },
+  "webgpu:compat,api,validation,render_pipeline,shader_module:unsupportedStorageTextureFormats,renderPipeline:*": { "subcaseMS": 0.601 },
   "webgpu:compat,api,validation,render_pipeline,vertex_state:maxVertexAttributesVertexIndexInstanceIndex:*": { "subcaseMS": 3.700 },
   "webgpu:compat,api,validation,texture,createTexture:depthOrArrayLayers_incompatible_with_textureBindingViewDimension:*": { "subcaseMS": 12.712 },
   "webgpu:compat,api,validation,texture,createTexture:format_reinterpretation:*": { "subcaseMS": 7.012 },
   "webgpu:compat,api,validation,texture,createTexture:invalidTextureBindingViewDimension:*": { "subcaseMS": 6.022 },
+  "webgpu:compat,api,validation,texture,createTexture:unsupportedStorageTextureFormats:*": { "subcaseMS": 0.601 },
   "webgpu:compat,api,validation,texture,createTexture:unsupportedTextureFormats:*": { "subcaseMS": 0.700 },
   "webgpu:compat,api,validation,texture,createTexture:unsupportedTextureViewFormats:*": { "subcaseMS": 0.601 },
   "webgpu:compat,api,validation,texture,cubeArray:cube_array:*": { "subcaseMS": 13.701 },
@@ -1640,6 +1645,7 @@
   "webgpu:shader,execution,robust_access:linear_memory:*": { "subcaseMS": 5.293 },
   "webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:*": { "subcaseMS": 6.487 },
   "webgpu:shader,execution,shader_io,compute_builtins:inputs:*": { "subcaseMS": 19.342 },
+  "webgpu:shader,execution,shader_io,fragment_builtins:inputs,front_facing:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,position:*": { "subcaseMS": 1.001 },
   "webgpu:shader,execution,shader_io,fragment_builtins:inputs,sample_index:*": { "subcaseMS": 1.001 },
diff --git a/src/webgpu/shader/execution/expression/call/builtin/atomics/atomicCompareExchangeWeak.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/atomics/atomicCompareExchangeWeak.spec.ts
index 2556bb744ba9..85cc5cce4dec 100644
--- a/src/webgpu/shader/execution/expression/call/builtin/atomics/atomicCompareExchangeWeak.spec.ts
+++ b/src/webgpu/shader/execution/expression/call/builtin/atomics/atomicCompareExchangeWeak.spec.ts
@@ -339,6 +339,11 @@ struct __atomic_compare_exchange_result<T> {
     const numInvocations = t.params.workgroupSize;
     const scalarType = t.params.scalarType;
 
+    t.skipIf(
+      numInvocations > t.device.limits.maxComputeWorkgroupSizeX,
+      `${numInvocations} > maxComputeWorkgroupSizeX(${t.device.limits.maxComputeWorkgroupSizeX})`
+    );
+
     // Number of times each workgroup attempts to exchange the same value to the same memory address
     const numWrites = 4;
 
@@ -556,6 +561,11 @@ struct __atomic_compare_exchange_result<T> {
     const numInvocations = t.params.workgroupSize;
     const scalarType = t.params.scalarType;
 
+    t.skipIf(
+      numInvocations > t.device.limits.maxComputeWorkgroupSizeX,
+      `${numInvocations} > maxComputeWorkgroupSizeX(${t.device.limits.maxComputeWorkgroupSizeX})`
+    );
+
     // Number of times each workgroup attempts to exchange the same value to the same memory address
     const numWrites = 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 76c9b85ff6d4..1cd09486c1b7 100644
--- a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts
+++ b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts
@@ -1,16 +1,16 @@
 export const description = `Test fragment shader builtin variables and inter-stage variables
 
 * test builtin(position)
-* test interpolation
+* test @interpolate
+* test builtin(sample_index)
+* test builtin(front_facing)
 
-The current tests draw a single triangle with clip space coordinates [1, 1], [-3, 1], [1, -3].
-This means they render to all pixels in the textures. To fully test centroid interpolation
-probably requires drawing various triangles that only cover certain samples. That is TBD.
+Note: @interpolate settings and sample_index affect whether or not the fragment shader
+is evaluated per-fragment or per-sample. With @interpolate(, sample) or usage of
+@builtin(sample_index) the fragment shader should be executed per-sample.
 
 TODO:
-* test sample interpolation (see MAINTENANCE_TODOs below)
 * test centroid interpolation (see MAINTENANCE_TODOs below)
-* test front_facing
 * test frag_depth
 `;
 
@@ -173,6 +173,8 @@ function copyRGBA8EncodedFloatTexturesToBufferIncludingMultisampledTextures(
 }
 
 /* column constants */
+const kX = 0;
+const kY = 1;
 const kZ = 2;
 const kW = 3;
 
@@ -263,7 +265,33 @@ function calcBarycentricCoordinates(trianglePoints: number[][], p: number[]) {
   return [u, v, w];
 }
 
+/**
+ * Returns true if point is inside triangle
+ */
+function isInsideTriangle(barycentricCoords: number[]) {
+  for (const v of barycentricCoords) {
+    if (v < 0 || v > 1) {
+      return false;
+    }
+  }
+  return true;
+}
+
+/**
+ * Returns true if windowPoints define a clockwise triangle
+ */
+function isTriangleClockwise(windowPoints: readonly number[][]) {
+  let sum = 0;
+  for (let i = 0; i < 3; ++i) {
+    const p0 = windowPoints[i];
+    const p1 = windowPoints[(i + 1) % 3];
+    sum += p0[kX] * p1[kY] - p1[kX] * p0[kY];
+  }
+  return sum >= 0;
+}
+
 type FragData = {
+  baseVertexIndex: number;
   fragmentPoint: readonly number[];
   fragmentBarycentricCoords: readonly number[];
   sampleBarycentricCoords: readonly number[];
@@ -271,6 +299,7 @@ type FragData = {
   ndcPoints: readonly number[][];
   windowPoints: readonly number[][];
   sampleIndex: number;
+  frontFacing: boolean;
 };
 
 /**
@@ -286,6 +315,7 @@ function generateFragmentInputs({
   height,
   nearFar,
   sampleCount,
+  frontFace,
   clipSpacePoints,
   interpolateFn,
 }: {
@@ -293,6 +323,7 @@ function generateFragmentInputs({
   height: number;
   nearFar: readonly number[];
   sampleCount: number;
+  frontFace?: GPUFrontFace;
   clipSpacePoints: readonly number[][];
   interpolateFn: (fragData: FragData) => number[];
 }) {
@@ -300,35 +331,49 @@ function generateFragmentInputs({
 
   const viewport = [0, 0, width, height, ...nearFar];
 
-  const ndcPoints = clipSpacePoints.map(clipSpaceToNDC);
-  const windowPoints = ndcPoints.map(p => ndcToWindow(p, viewport));
-  const windowPoints2D = windowPoints.map(p => p.slice(0, 2));
-
-  const fragmentOffsets = getMultisampleFragmentOffsets(sampleCount)!;
-  for (let y = 0; y < height; ++y) {
-    for (let x = 0; x < width; ++x) {
-      for (let s = 0; s < sampleCount; ++s) {
-        const fragmentPoint = [x + 0.5, y + 0.5];
-        const multisampleOffset = fragmentOffsets[s];
-        const sampleFragmentPoint = [x + multisampleOffset[0], y + multisampleOffset[1]];
-        const fragmentBarycentricCoords = calcBarycentricCoordinates(windowPoints2D, fragmentPoint);
-        const sampleBarycentricCoords = calcBarycentricCoordinates(
-          windowPoints2D,
-          sampleFragmentPoint
-        );
-
-        const output = interpolateFn({
-          fragmentPoint,
-          fragmentBarycentricCoords,
-          sampleBarycentricCoords,
-          clipSpacePoints,
-          ndcPoints,
-          windowPoints,
-          sampleIndex: s,
-        });
-
-        const offset = ((y * width + x) * sampleCount + s) * 4;
-        expected.set(output, offset);
+  // For each triangle
+  for (let vertexIndex = 0; vertexIndex < clipSpacePoints.length; vertexIndex += 3) {
+    const ndcPoints = clipSpacePoints.slice(vertexIndex, vertexIndex + 3).map(clipSpaceToNDC);
+    const windowPoints = ndcPoints.map(p => ndcToWindow(p, viewport));
+    const windowPoints2D = windowPoints.map(p => p.slice(0, 2));
+
+    const cw = isTriangleClockwise(windowPoints2D);
+    const frontFacing = frontFace === 'cw' ? cw : !cw;
+    const fragmentOffsets = getMultisampleFragmentOffsets(sampleCount)!;
+
+    for (let y = 0; y < height; ++y) {
+      for (let x = 0; x < width; ++x) {
+        for (let sampleIndex = 0; sampleIndex < sampleCount; ++sampleIndex) {
+          const fragmentPoint = [x + 0.5, y + 0.5];
+          const multisampleOffset = fragmentOffsets[sampleIndex];
+          const sampleFragmentPoint = [x + multisampleOffset[0], y + multisampleOffset[1]];
+          const fragmentBarycentricCoords = calcBarycentricCoordinates(
+            windowPoints2D,
+            fragmentPoint
+          );
+          const sampleBarycentricCoords = calcBarycentricCoordinates(
+            windowPoints2D,
+            sampleFragmentPoint
+          );
+
+          const inside = isInsideTriangle(sampleBarycentricCoords);
+          if (inside) {
+            const output = interpolateFn({
+              baseVertexIndex: vertexIndex,
+              fragmentPoint,
+              fragmentBarycentricCoords,
+              sampleBarycentricCoords,
+              clipSpacePoints,
+              ndcPoints,
+              windowPoints,
+              sampleIndex,
+              frontFacing,
+            });
+
+            const offset = ((y * width + x) * sampleCount + sampleIndex) * 4;
+            expected.set(output, offset);
+          }
+        }
       }
     }
   }
@@ -366,29 +411,31 @@ function createInterStageInterpolationFn(
   sampling: InterpolationSampling | undefined
 ) {
   return function ({
+    baseVertexIndex,
     fragmentBarycentricCoords,
     sampleBarycentricCoords,
     clipSpacePoints,
   }: FragData) {
+    const triangleInterStagePoints = interStagePoints.slice(baseVertexIndex, baseVertexIndex + 3);
     const barycentricCoords =
       sampling === 'center' ? fragmentBarycentricCoords : sampleBarycentricCoords;
     switch (type) {
       case 'perspective':
-        return interStagePoints[0].map((_, colNum: number) =>
+        return triangleInterStagePoints[0].map((_, colNum: number) =>
           perspectiveInterpolation(
             barycentricCoords,
             clipSpacePoints,
-            getColumn(interStagePoints, colNum)
+            getColumn(triangleInterStagePoints, colNum)
           )
         );
         break;
       case 'linear':
-        return interStagePoints[0].map((_, colNum: number) =>
-          linearInterpolation(barycentricCoords, getColumn(interStagePoints, colNum))
+        return triangleInterStagePoints[0].map((_, colNum: number) =>
+          linearInterpolation(barycentricCoords, getColumn(triangleInterStagePoints, colNum))
         );
         break;
       case 'flat':
-        return interStagePoints[0];
+        return triangleInterStagePoints[0];
         break;
       default:
         unreachable();
@@ -403,6 +450,13 @@ function computeFragmentSampleIndex({ sampleIndex }: FragData) {
   return [sampleIndex, 0, 0, 0];
 }
 
+/**
+ * Computes 'builtin(front_facing)'
+ */
+function computeFragmentFrontFacing({ frontFacing }: FragData) {
+  return [frontFacing ? 1 : 0, 0, 0, 0];
+}
+
 /**
  * Renders float32 fragment shader inputs values to 4 rgba8unorm textures that
  * can be multisampled textures. It stores each of the channels, r, g, b, a of
@@ -430,8 +484,10 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
     width,
     height,
     nearFar,
+    frontFace,
     clipSpacePoints,
     interStagePoints,
+    fragInCode,
     outputCode,
   }: {
     interpolationType: InterpolationType;
@@ -439,9 +495,11 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
     width: number;
     height: number;
     sampleCount: number;
+    frontFace?: GPUFrontFace;
     nearFar: readonly number[];
     clipSpacePoints: readonly number[][];
     interStagePoints: readonly number[][];
+    fragInCode: string;
     outputCode: string;
   }
 ) {
@@ -461,12 +519,6 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
         @location(0) @interpolate(${interpolate}) interpolatedValue: vec4f,
       };
 
-      struct VertexIn {
-        @builtin(position) position: vec4f,
-        @location(0) @interpolate(${interpolate}) interpolatedValue: vec4f,
-        @builtin(sample_index) sampleIndex: u32,
-      };
-
       @vertex fn vs(@builtin(vertex_index) vNdx: u32) -> VertexOut {
         let pos = array(
           ${clipSpacePoints.map(p => `vec4f(${p.join(', ')})`).join(', ')}
@@ -481,6 +533,12 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
         return v;
       }
 
+      struct FragmentIn {
+        @builtin(position) position: vec4f,
+        @location(0) @interpolate(${interpolate}) interpolatedValue: vec4f,
+        ${fragInCode}
+      };
+
       struct FragOut {
         @location(0) out0: vec4f,
         @location(1) out1: vec4f,
@@ -497,7 +555,7 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
         );
       }
 
-      @fragment fn fs(vin: VertexIn) -> FragOut {
+      @fragment fn fs(fin: FragmentIn) -> FragOut {
         var f: FragOut;
         let v = ${outputCode};
         let u = bitcast<vec4u>(v);
@@ -505,7 +563,7 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
         f.out1 = u32ToRGBAUnorm(u[1]);
         f.out2 = u32ToRGBAUnorm(u[2]);
         f.out3 = u32ToRGBAUnorm(u[3]);
-        _ = vin.interpolatedValue;
+        _ = fin.interpolatedValue;
         return f;
       }
     `,
@@ -536,6 +594,11 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues(
       entryPoint: 'fs',
       targets: textures.map(() => ({ format: 'rgba8unorm' })),
     },
+    ...(frontFace && {
+      primitive: {
+        frontFace,
+      },
+    }),
     multisample: {
       count: sampleCount,
     },
@@ -581,17 +644,17 @@ function checkSampleRectsApproximatelyEqual({
   sampleCount,
   actual,
   expected,
-  maxFractionalDiff,
+  maxDiffULPsForFloatFormat,
 }: {
   width: number;
   height: number;
   sampleCount: number;
   actual: Float32Array;
   expected: Float32Array;
-  maxFractionalDiff: number;
+  maxDiffULPsForFloatFormat: number;
 }) {
   const subrectOrigin = [0, 0, 0];
-  const subrectSize = [width, height, 1];
+  const subrectSize = [width * sampleCount, height, 1];
   const areaDesc = {
     bytesPerRow: width * sampleCount * 4 * 4,
     rowsPerImage: height,
@@ -614,9 +677,9 @@ function checkSampleRectsApproximatelyEqual({
   const failedPixelsMessage = findFailedPixels(
     format,
     { x: 0, y: 0, z: 0 },
-    { width, height, depthOrArrayLayers: 1 },
+    { width: width * sampleCount, height, depthOrArrayLayers: 1 },
     { actTexelView, expTexelView },
-    { maxFractionalDiff }
+    { maxDiffULPsForFloatFormat }
   );
 
   if (failedPixelsMessage !== undefined) {
@@ -634,6 +697,8 @@ g.test('inputs,position')
   .desc(
     `
     Test fragment shader builtin(position) values.
+
+    Note: @builtin(position) is always a fragment position, never a sample position.
   `
   )
   .params(u =>
@@ -642,13 +707,11 @@ g.test('inputs,position')
       .combine('sampleCount', [1, 4] as const)
       .combine('interpolation', [
         { type: 'perspective', sampling: 'center' },
-        // MAINTENANCE_TODO: enable these tests.
-        // { type: 'perspective', sampling: 'centroid' },
-        // { type: 'perspective', sampling: 'sample' },
+        { type: 'perspective', sampling: 'centroid' },
+        { type: 'perspective', sampling: 'sample' },
         { type: 'linear', sampling: 'center' },
-        // MAINTENANCE_TODO: enable these tests.
-        // { type: 'linear', sampling: 'centroid' },
-        // { type: 'linear', sampling: 'sample' },
+        { type: 'linear', sampling: 'centroid' },
+        { type: 'linear', sampling: 'sample' },
         { type: 'flat' },
       ] as const)
   )
@@ -688,7 +751,8 @@ g.test('inputs,position')
       nearFar,
       clipSpacePoints,
       interStagePoints,
-      outputCode: 'vin.position',
+      fragInCode: '',
+      outputCode: 'fin.position',
     });
 
     const expected = generateFragmentInputs({
@@ -700,6 +764,12 @@ g.test('inputs,position')
       interpolateFn: computeFragmentPosition,
     });
 
+    // Since @builtin(position) is always a fragment position, never a sample position, check
+    // the first coordinate. It should be 0.5, 0.5 always. This is just to double check
+    // that computeFragmentPosition is generating the correct values.
+    assert(expected[0] === 0.5);
+    assert(expected[1] === 0.5);
+
     t.expectOK(
       checkSampleRectsApproximatelyEqual({
         width,
@@ -707,7 +777,7 @@ g.test('inputs,position')
         sampleCount,
         actual,
         expected,
-        maxFractionalDiff: 0.000001,
+        maxDiffULPsForFloatFormat: 2,
       })
     );
   });
@@ -726,11 +796,11 @@ g.test('inputs,interStage')
         { type: 'perspective', sampling: 'center' },
         // MAINTENANCE_TODO: enable these tests.
         // { type: 'perspective', sampling: 'centroid' },
-        // { type: 'perspective', sampling: 'sample' },
+        { type: 'perspective', sampling: 'sample' },
         { type: 'linear', sampling: 'center' },
         // MAINTENANCE_TODO: enable these tests.
         // { type: 'linear', sampling: 'centroid' },
-        // { type: 'linear', sampling: 'sample' },
+        { type: 'linear', sampling: 'sample' },
         { type: 'flat' },
       ] as const)
   )
@@ -770,7 +840,8 @@ g.test('inputs,interStage')
       nearFar,
       clipSpacePoints,
       interStagePoints,
-      outputCode: 'vin.interpolatedValue',
+      fragInCode: '',
+      outputCode: 'fin.interpolatedValue',
     });
 
     const expected = generateFragmentInputs({
@@ -789,7 +860,7 @@ g.test('inputs,interStage')
         sampleCount,
         actual,
         expected,
-        maxFractionalDiff: 0.00001,
+        maxDiffULPsForFloatFormat: 3,
       })
     );
   });
@@ -850,7 +921,8 @@ g.test('inputs,sample_index')
       nearFar,
       clipSpacePoints,
       interStagePoints,
-      outputCode: 'vec4f(f32(vin.sampleIndex), 0, 0, 0)',
+      fragInCode: `@builtin(sample_index) sampleIndex: u32,`,
+      outputCode: 'vec4f(f32(fin.sampleIndex), 0, 0, 0)',
     });
 
     const expected = generateFragmentInputs({
@@ -869,7 +941,124 @@ g.test('inputs,sample_index')
         sampleCount,
         actual,
         expected,
-        maxFractionalDiff: 0.00001,
+        maxDiffULPsForFloatFormat: 1,
+      })
+    );
+  });
+
+g.test('inputs,front_facing')
+  .desc(
+    `
+    Test fragment shader builtin(front_facing) values.
+
+    Draws a quad from 2 triangles that entirely cover clip space. (see diagram below in code)
+    One triangle is clockwise, the other is counter clockwise. The triangles
+    bisect pixels so that different samples are covered by each triangle so that some
+    samples should get different values for front_facing for the same fragment.
+  `
+  )
+  .params(u =>
+    u //
+      .combine('nearFar', [[0, 1] as const, [0.25, 0.75] as const] as const)
+      .combine('sampleCount', [1, 4] as const)
+      .combine('frontFace', ['cw', 'ccw'] as const)
+      .combine('interpolation', [
+        { type: 'perspective', sampling: 'center' },
+        { type: 'perspective', sampling: 'centroid' },
+        { type: 'perspective', sampling: 'sample' },
+        { type: 'linear', sampling: 'center' },
+        { type: 'linear', sampling: 'centroid' },
+        { type: 'linear', sampling: 'sample' },
+        { type: 'flat' },
+      ] as const)
+  )
+  .beforeAllSubcases(t => {
+    const {
+      interpolation: { type, sampling },
+    } = t.params;
+    t.skipIfInterpolationTypeOrSamplingNotSupported({ type, sampling });
+  })
+  .fn(async t => {
+    const {
+      nearFar,
+      sampleCount,
+      frontFace,
+      interpolation: { type, sampling },
+    } = t.params;
+    //
+    // We're drawing 2 triangles starting at y = -2 to y = +2
+    //
+    //  -1   0   1
+    //   +===+===+  2
+    //   |\  |   |
+    //   +---+---+  1  <---
+    //   |  \|   |       |
+    //   +---+---+  0    | viewport
+    //   |   |\  |       |
+    //   +---+---+ -1  <---
+    //   |   |  \|
+    //   +===+===+ -2
+
+    // prettier-ignore
+    const clipSpacePoints = [
+      // ccw
+      [-1, -2, 0, 1],
+      [ 1, -2, 0, 1],
+      [-1,  2, 0, 1],
+
+      // cw
+      [ 1, -2, 0, 1],
+      [-1,  2, 0, 1],
+      [ 1,  2, 0, 1],
+    ];
+
+    const interStagePoints = [
+      [1, 2, 3, 4],
+      [5, 6, 7, 8],
+      [9, 10, 11, 12],
+
+      [13, 14, 15, 16],
+      [17, 18, 19, 20],
+      [21, 22, 23, 24],
+    ];
+
+    const width = 4;
+    const height = 4;
+    const actual = await renderFragmentShaderInputsTo4TexturesAndReadbackValues(t, {
+      interpolationType: type,
+      interpolationSampling: sampling,
+      frontFace,
+      sampleCount,
+      width,
+      height,
+      nearFar,
+      clipSpacePoints,
+      interStagePoints,
+      fragInCode: '@builtin(front_facing) frontFacing: bool,',
+      outputCode: 'vec4f(select(0.0, 1.0, fin.frontFacing), 0, 0, 0)',
+    });
+
+    const expected = generateFragmentInputs({
+      width,
+      height,
+      nearFar,
+      sampleCount,
+      clipSpacePoints,
+      frontFace,
+      interpolateFn: computeFragmentFrontFacing,
+    });
+
+    // Double check, first corner should be different than last based on the triangles we are drawing.
+    assert(expected[0] !== expected[expected.length - 4]);
+
+    t.expectOK(
+      checkSampleRectsApproximatelyEqual({
+        width,
+        height,
+        sampleCount,
+        actual,
+        expected,
+        maxDiffULPsForFloatFormat: 0,
       })
     );
   });