chore: Update WebGPU CTS (#33990)

* chore: Update WebGPU CTS

d473d09475
Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* Update expectations

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

* Allow multiple process testing for CTS

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

---------

Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com>

python/servo/try_parser.py

@@ -102,7 +102,7 @@ def handle_preset(s: str) -> Optional[JobConfig]:
     elif s == "webgpu":
         return JobConfig("WebGPU CTS", Workflow.LINUX,
                          wpt_layout=Layout.layout2020,  # reftests are mode for new layout
-                         wpt_args="--processes 1 _webgpu",  # run only webgpu cts
+                         wpt_args="_webgpu",  # run only webgpu cts
                          profile="production",  # WebGPU works to slow with debug assert
                          unit_tests=False)  # production profile does not work with unit-tests
     elif s in ["lint", "tidy"]:

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_colorspace_bgra8unorm.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_colorspace_bgra8unorm.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_colorspace_rgba16float.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_colorspace_rgba16float.https.html]
   expected:
-    if os == "linux" and not debug: TIMEOUT
+    if os == "linux" and not debug: FAIL

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_colorspace_rgba8unorm.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_colorspace_rgba8unorm.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_bgra8unorm_copy.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_complex_bgra8unorm_copy.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_bgra8unorm_draw.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_complex_bgra8unorm_draw.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_rgba16float_copy.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_complex_rgba16float_copy.https.html]
   expected:
-    if os == "linux" and not debug: TIMEOUT
+    if os == "linux" and not debug: FAIL

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_rgba16float_draw.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_complex_rgba16float_draw.https.html]
   expected:
-    if os == "linux" and not debug: TIMEOUT
+    if os == "linux" and not debug: FAIL

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_rgba8unorm_copy.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_complex_rgba8unorm_copy.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_complex_rgba8unorm_draw.https.html.ini

@@ -1,2 +1,6 @@
 [canvas_complex_rgba8unorm_draw.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: FAIL
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_composite_alpha_bgra8unorm_premultiplied_copy.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_composite_alpha_bgra8unorm_premultiplied_copy.https.html]
   expected:
-    if os == "linux" and not debug: PASS
+    if os == "linux" and not debug: [PASS, FAIL]

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_composite_alpha_bgra8unorm_premultiplied_draw.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_composite_alpha_bgra8unorm_premultiplied_draw.https.html]
   expected:
-    if os == "linux" and not debug: PASS
+    if os == "linux" and not debug: [PASS, FAIL]

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_composite_alpha_rgba8unorm_opaque_copy.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_composite_alpha_rgba8unorm_opaque_copy.https.html]
   expected:
-    if os == "linux" and not debug: PASS
+    if os == "linux" and not debug: [PASS, FAIL]

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_composite_alpha_rgba8unorm_premultiplied_copy.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_composite_alpha_rgba8unorm_premultiplied_copy.https.html]
   expected:
-    if os == "linux" and not debug: PASS
+    if os == "linux" and not debug: [PASS, FAIL]

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/canvas_composite_alpha_rgba8unorm_premultiplied_draw.https.html.ini

@@ -1,3 +1,3 @@
 [canvas_composite_alpha_rgba8unorm_premultiplied_draw.https.html]
   expected:
-    if os == "linux" and not debug: PASS
+    if os == "linux" and not debug: [CRASH, PASS]

tests/wpt/webgpu/meta/webgpu/webgpu/web_platform/reftests/resize_observer.https.html.ini

@@ -1,2 +1,6 @@
 [resize_observer.https.html]
-  expected: TIMEOUT
+  expected:
+    if os == "win": TIMEOUT
+    if os == "linux" and debug: TIMEOUT
+    if os == "linux" and not debug: PASS
+    if os == "mac": TIMEOUT

tests/wpt/webgpu/tests/checkout_commit.txt

@@ -1 +1 @@
-50b6e7a7435e8d1a973cbf67347938ce05188df0
+d473d09475bffec9569fe5c45834bb6aaad44818

tests/wpt/webgpu/tests/webgpu/common/framework/test_config.js

@@ -25,6 +25,18 @@
 
 
 
+
+
+
+
+
+
+
+
+
+
+
+
 
 
 
@@ -40,7 +52,8 @@
 
 export const globalTestConfig = {
   enableDebugLogs: false,
-  maxSubcasesInFlight: 500,
+  maxSubcasesInFlight: 100,
+  subcasesBetweenAttemptingGC: 5000,
   testHeartbeatCallback: () => {},
   noRaceWithRejectOnTimeout: false,
   unrollConstEvalLoops: false,

tests/wpt/webgpu/tests/webgpu/common/internal/test_group.js

@@ -31,6 +31,7 @@ import {
   stringifyPublicParamsUniquely } from
 '../internal/query/stringify_params.js';
 import { validQueryPart } from '../internal/query/validQueryPart.js';
+import { attemptGarbageCollection } from '../util/collect_garbage.js';
 
 import { assert, unreachable } from '../util/util.js';
 
@@ -620,7 +621,7 @@ class RunCaseSpecific {
             const subcasePrefix = 'subcase: ' + stringifyPublicParams(subParams);
             const subRec = new Proxy(rec, {
               get: (target, k) => {
-                const prop = TestCaseRecorder.prototype[k];
+                const prop = rec[k] ?? TestCaseRecorder.prototype[k];
                 if (typeof prop === 'function') {
                   testHeartbeatCallback();
                   return function (...args) {
@@ -696,6 +697,7 @@ class RunCaseSpecific {
                 subRec.threw(ex);
               }
             }).
+            finally(attemptGarbageCollectionIfDue).
             finally(subcaseFinishedCallback);
 
             allPreviousSubcasesFinalizedPromise = allPreviousSubcasesFinalizedPromise.then(
@@ -711,13 +713,17 @@ class RunCaseSpecific {
             rec.skipped(new SkipTestCase('all subcases were skipped'));
           }
         } else {
-          await this.runTest(
-            rec,
-            sharedState,
-            this.params,
-            /* throwSkip */false,
-            getExpectedStatus(selfQuery)
-          );
+          try {
+            await this.runTest(
+              rec,
+              sharedState,
+              this.params,
+              /* throwSkip */false,
+              getExpectedStatus(selfQuery)
+            );
+          } finally {
+            await attemptGarbageCollectionIfDue();
+          }
         }
       } finally {
         testHeartbeatCallback();
@@ -742,4 +748,29 @@ class RunCaseSpecific {
       logToWebSocket(JSON.stringify(msg));
     }
   }
-}
+}
+
+
+
+
+
+
+
+
+
+
+
+
+/** Every `subcasesBetweenAttemptingGC` calls to this function will `attemptGarbageCollection()`. */
+const attemptGarbageCollectionIfDue = (() => {
+  // This state is global because garbage is global.
+  let subcasesSinceLastGC = 0;
+
+  return async function attemptGarbageCollectionIfDue() {
+    subcasesSinceLastGC++;
+    if (subcasesSinceLastGC >= globalTestConfig.subcasesBetweenAttemptingGC) {
+      subcasesSinceLastGC = 0;
+      return attemptGarbageCollection();
+    }
+  };
+})();

tests/wpt/webgpu/tests/webgpu/common/internal/version.js

@@ -1,3 +1,3 @@
 // AUTO-GENERATED - DO NOT EDIT. See tools/gen_version.
 
-export const version = '50b6e7a7435e8d1a973cbf67347938ce05188df0';
+export const version = 'd473d09475bffec9569fe5c45834bb6aaad44818';

tests/wpt/webgpu/tests/webgpu/common/util/navigator_gpu.js

@@ -68,12 +68,11 @@ export function getGPU(recorder) {
     {
       const promise = oldFn.call(this, { ...defaultRequestAdapterOptions, ...options });
       if (recorder) {
-        void promise.then(async (adapter) => {
+        void promise.then((adapter) => {
           if (adapter) {
-            // MAINTENANCE_TODO: Remove requestAdapterInfo when info is implemented.
-            const info = adapter.info || (await adapter.requestAdapterInfo());
-            const infoString = `Adapter: ${info.vendor} / ${info.architecture} / ${info.device}`;
-            recorder.debug(new ErrorWithExtra(infoString, () => ({ adapterInfo: info })));
+            const adapterInfo = adapter.info;
+            const infoString = `Adapter: ${adapterInfo.vendor} / ${adapterInfo.architecture} / ${adapterInfo.device}`;
+            recorder.debug(new ErrorWithExtra(infoString, () => ({ adapterInfo })));
           }
         });
       }

tests/wpt/webgpu/tests/webgpu/cts.https.html

@@ -306,6 +306,8 @@
 <meta name=variant content='?q=webgpu:api,validation,buffer,mapping:unmap,state,mappingPending:*'>
 <meta name=variant content='?q=webgpu:api,validation,buffer,mapping:gc_behavior,mappedAtCreation:*'>
 <meta name=variant content='?q=webgpu:api,validation,buffer,mapping:gc_behavior,mapAsync:*'>
+<meta name=variant content='?q=webgpu:api,validation,capability_checks,features,clip_distances:createRenderPipeline,at_over:*'>
+<meta name=variant content='?q=webgpu:api,validation,capability_checks,features,clip_distances:createRenderPipeline,max_vertex_output_location:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,features,query_types:createQuerySet:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,features,query_types:timestamp:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,features,texture_formats:texture_descriptor:*'>
@@ -346,7 +348,6 @@
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxComputeWorkgroupsPerDimension:validate:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxDynamicStorageBuffersPerPipelineLayout:createBindGroupLayout,at_over:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxDynamicUniformBuffersPerPipelineLayout:createBindGroupLayout,at_over:*'>
-<meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxInterStageShaderComponents:createRenderPipeline,at_over:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxInterStageShaderVariables:createRenderPipeline,at_over:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxSampledTexturesPerShaderStage:createBindGroupLayout,at_over:*'>
 <meta name=variant content='?q=webgpu:api,validation,capability_checks,limits,maxSampledTexturesPerShaderStage:createPipelineLayout,at_over:*'>
@@ -467,6 +468,7 @@
 <meta name=variant content='?q=webgpu:api,validation,createView:mip_levels:*'>
 <meta name=variant content='?q=webgpu:api,validation,createView:cube_faces_square:*'>
 <meta name=variant content='?q=webgpu:api,validation,createView:texture_state:*'>
+<meta name=variant content='?q=webgpu:api,validation,createView:texture_view_usage:*'>
 <meta name=variant content='?q=webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,command_encoder:*'>
 <meta name=variant content='?q=webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,render_compute_pass:*'>
 <meta name=variant content='?q=webgpu:api,validation,encoding,beginComputePass:timestampWrites,query_set_type:*'>
@@ -749,6 +751,7 @@
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,depth_stencil_state:depth_bias:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,depth_stencil_state:stencil_test:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,depth_stencil_state:stencil_write:*'>
+<meta name=variant content='?q=webgpu:api,validation,render_pipeline,float32_blendable:create_render_pipeline:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:color_target_exists:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:targets_format_is_color_format:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:targets_format_renderable:*'>
@@ -760,6 +763,8 @@
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:targets_write_mask:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:pipeline_output_targets:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:pipeline_output_targets,blend:*'>
+<meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:dual_source_blending,color_target_count:*'>
+<meta name=variant content='?q=webgpu:api,validation,render_pipeline,fragment_state:dual_source_blending,use_blend_src:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:location,mismatch:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:location,superset:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:location,subset:*'>
@@ -767,8 +772,8 @@
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:interpolation_type:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:interpolation_sampling:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:max_shader_variable_location:*'>
-<meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:max_components_count,output:*'>
-<meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:max_components_count,input:*'>
+<meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:max_variables_count,output:*'>
+<meta name=variant content='?q=webgpu:api,validation,render_pipeline,inter_stage:max_variables_count,input:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,misc:basic:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,misc:no_attachment:*'>
 <meta name=variant content='?q=webgpu:api,validation,render_pipeline,misc:vertex_state_only:*'>
@@ -834,6 +839,7 @@
 <meta name=variant content='?q=webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,set_bind_group_on_same_index_depth_stencil_texture:*'>
 <meta name=variant content='?q=webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,set_unused_bind_group:*'>
 <meta name=variant content='?q=webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,texture_usages_in_copy_and_render_pass:*'>
+<meta name=variant content='?q=webgpu:api,validation,resource_usages,texture,in_render_misc:subresources,texture_view_usages:*'>
 <meta name=variant content='?q=webgpu:api,validation,shader_module,entry_point:compute:*'>
 <meta name=variant content='?q=webgpu:api,validation,shader_module,entry_point:vertex:*'>
 <meta name=variant content='?q=webgpu:api,validation,shader_module,entry_point:fragment:*'>
@@ -900,12 +906,14 @@
 <meta name=variant content='?q=webgpu:compat,api,validation,encoding,programmable,pipeline_bind_group_compat:twoDifferentTextureViews,compute_pass,unused:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,depth_stencil_state:depthBiasClamp:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,fragment_state:colorState:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:sample_mask:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:sample_index:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:interpolate:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:unsupportedStorageTextureFormats,computePipeline:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:unsupportedStorageTextureFormats,renderPipeline:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:textureLoad_with_depth_textures,computePipeline:*'>
+<meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:textureLoad_with_depth_textures,renderPipeline:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,render_pipeline,vertex_state:maxVertexAttributesVertexIndexInstanceIndex:*'>
-<meta name=variant content='?q=webgpu:compat,api,validation,shader_module,shader_module:sample_mask:*'>
-<meta name=variant content='?q=webgpu:compat,api,validation,shader_module,shader_module:sample_index:*'>
-<meta name=variant content='?q=webgpu:compat,api,validation,shader_module,shader_module:interpolate:*'>
-<meta name=variant content='?q=webgpu:compat,api,validation,shader_module,shader_module:unsupportedStorageTextureFormats:*'>
-<meta name=variant content='?q=webgpu:compat,api,validation,shader_module,shader_module:textureLoad_with_depth_textures:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,texture,createTexture:unsupportedTextureFormats:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,texture,createTexture:unsupportedTextureViewFormats:*'>
 <meta name=variant content='?q=webgpu:compat,api,validation,texture,createTexture:invalidTextureBindingViewDimension:*'>
@@ -950,6 +958,7 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,matrix,index:concrete_float_element:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,matrix,index:abstract_float_column:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,matrix,index:abstract_float_element:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,access,matrix,index:non_const_index:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,structure,index:buffer:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,structure,index:buffer_align:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,access,structure,index:buffer_size:*'>
@@ -1508,6 +1517,16 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,pow:abstract_float:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,pow:f32:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,pow:f16:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadBroadcast:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadBroadcast:compute,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadBroadcast:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadBroadcast:fragment,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadBroadcast:fragment,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadSwap:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadSwap:compute,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadSwap:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadSwap:fragment,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quadSwap:fragment,split:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,quantizeToF16:f32:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,radians:abstract_float:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,radians:f32:*'>
@@ -1562,6 +1581,35 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,step:f16:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,storageBarrier:stage:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,storageBarrier:barrier:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAdd:fp_accuracy:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAdd:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAdd:fragment:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAdd:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAll:compute,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAll:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAll:fragment,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAll:fragment,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAny:compute,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAny:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAny:fragment,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupAny:fragment,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBallot:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBallot:fragment,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBallot:fragment:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBallot:predicate:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBallot:predicate_and_control_flow:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBitwise:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBitwise:compute,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBitwise:compute,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBitwise:fragment,all_active:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBitwise:fragment,split:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:workgroup_uniform_load:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupBroadcast:fragment:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupMul:fp_accuracy:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupMul:data_types:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupMul:fragment:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,subgroupMul:compute,split:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,tan:abstract_float:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,tan:f32:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,tan:f16:*'>
@@ -1582,8 +1630,8 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGather:depth_array_3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:array_2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:array_3d_coords:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_array_2d_coords:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_array_3d_coords:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_2d_coords:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_1d:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_2d:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_3d:*'>
@@ -1591,7 +1639,10 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:depth:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:external:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:arrayed:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:storage_texel_formats:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_1d:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_2d:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_2d_array:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_3d:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureNumLayers:sampled:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureNumLayers:arrayed:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureNumLayers:storage:*'>
@@ -1601,7 +1652,6 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureNumSamples:depth:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:sampled_1d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:sampled_2d_coords:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:sampled_2d_coords,derivatives:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:sampled_3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:depth_2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:sampled_array_2d_coords:*'>
@@ -1609,6 +1659,7 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:depth_3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:depth_array_2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSample:depth_array_3d_coords:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleBaseClampToEdge:2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleBias:sampled_2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleBias:sampled_3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleBias:arrayed_2d_coords:*'>
@@ -1617,8 +1668,6 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompare:3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompare:arrayed_2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompare:arrayed_3d_coords:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompareLevel:stage:*'>
-<meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompareLevel:control_flow:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompareLevel:2d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompareLevel:3d_coords:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,call,builtin,textureSampleCompareLevel:arrayed_2d_coords:*'>
@@ -1693,6 +1742,7 @@
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,address_of_and_indirection:deref:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_index:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_member:*'>
+<meta name=variant content='?q=webgpu:shader,execution,expression,unary,address_of_and_indirection:deref_swizzle:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,af_arithmetic:negation:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,af_assignment:abstract:*'>
 <meta name=variant content='?q=webgpu:shader,execution,expression,unary,af_assignment:f32:*'>
@@ -1873,6 +1923,8 @@
 <meta name=variant content='?q=webgpu:shader,execution,padding:struct_explicit:*'>
 <meta name=variant content='?q=webgpu:shader,execution,padding:struct_nested:*'>
 <meta name=variant content='?q=webgpu:shader,execution,padding:array_of_vec3:*'>
+<meta name=variant content='?q=webgpu:shader,execution,padding:array_of_vec3h:*'>
+<meta name=variant content='?q=webgpu:shader,execution,padding:array_of_vec3h,elementwise:*'>
 <meta name=variant content='?q=webgpu:shader,execution,padding:array_of_struct:*'>
 <meta name=variant content='?q=webgpu:shader,execution,padding:vec3:*'>
 <meta name=variant content='?q=webgpu:shader,execution,padding:matCx3:*'>
@@ -1880,16 +1932,21 @@
 <meta name=variant content='?q=webgpu:shader,execution,robust_access:linear_memory:*'>
 <meta name=variant content='?q=webgpu:shader,execution,robust_access_vertex:vertex_buffer_access:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,compute_builtins:inputs:*'>
+<meta name=variant content='?q=webgpu:shader,execution,shader_io,compute_builtins:subgroup_size:*'>
+<meta name=variant content='?q=webgpu:shader,execution,shader_io,compute_builtins:subgroup_invocation_id:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,position:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,interStage,centroid:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,sample_index:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,front_facing:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:inputs,sample_mask:*'>
+<meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:subgroup_size:*'>
+<meta name=variant content='?q=webgpu:shader,execution,shader_io,fragment_builtins:subgroup_invocation_id:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,shared_structs:shared_with_buffer:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,shared_structs:shared_between_stages:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,shared_structs:shared_with_non_entry_point_function:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,user_io:passthrough:*'>
+<meta name=variant content='?q=webgpu:shader,execution,shader_io,vertex_builtins:outputs,clip_distances:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shader_io,workgroup_size:workgroup_size:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shadow:declaration:*'>
 <meta name=variant content='?q=webgpu:shader,execution,shadow:builtin:*'>
@@ -1925,6 +1982,7 @@
 <meta name=variant content='?q=webgpu:shader,execution,statement,increment_decrement:frexp_exp_increment:*'>
 <meta name=variant content='?q=webgpu:shader,execution,statement,increment_decrement:single_eval_increment:*'>
 <meta name=variant content='?q=webgpu:shader,execution,statement,increment_decrement:single_eval_decrement:*'>
+<meta name=variant content='?q=webgpu:shader,execution,statement,phony:executes:*'>
 <meta name=variant content='?q=webgpu:shader,execution,value_init:scalars:*'>
 <meta name=variant content='?q=webgpu:shader,execution,value_init:vec:*'>
 <meta name=variant content='?q=webgpu:shader,execution,value_init:mat:*'>
@@ -2024,6 +2082,11 @@
 <meta name=variant content='?q=webgpu:shader,validation,expression,binary,div_rem:scalar_vector_out_of_range:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,binary,div_rem:invalid_type_with_itself:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,binary,parse:all:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,binary,short_circuiting_and_or:scalar_vector:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_types:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_rhs_const:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_rhs_override:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,binary,short_circuiting_and_or:invalid_array_count_on_rhs:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,abs:values:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,abs:parameters:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,acos:values:*'>
@@ -2245,6 +2308,22 @@
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,pow:invalid_argument:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,pow:args:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,pow:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:id_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:id_constness:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadBroadcast:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quadSwap:stage:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quantizeToF16:values:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quantizeToF16:args:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,quantizeToF16:must_use:*'>
@@ -2297,6 +2376,73 @@
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,step:values:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,step:args:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,step:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAdd:invalid_types:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBallot:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBitwise:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:id_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:id_constness:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupElect:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMinMax:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:stage:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupMul:invalid_types:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:requires_subgroups:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:requires_subgroups_f16:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:early_eval:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:must_use:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:data_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:return_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:param2_type:*'>
+<meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,subgroupShuffle:stage:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,tan:values:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,tan:args:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,call,builtin,tan:must_use:*'>
@@ -2506,6 +2652,7 @@
 <meta name=variant content='?q=webgpu:shader,validation,expression,unary,logical_negation:scalar_vector:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,unary,logical_negation:invalid_types:*'>
 <meta name=variant content='?q=webgpu:shader,validation,expression,unary,logical_negation:parse:*'>
+<meta name=variant content='?q=webgpu:shader,validation,extension,clip_distances:use_clip_distances_requires_extension_enabled:*'>
 <meta name=variant content='?q=webgpu:shader,validation,extension,dual_source_blending:use_blend_src_requires_extension_enabled:*'>
 <meta name=variant content='?q=webgpu:shader,validation,extension,dual_source_blending:blend_src_syntax_validation:*'>
 <meta name=variant content='?q=webgpu:shader,validation,extension,dual_source_blending:blend_src_stage_input_output:*'>
@@ -2523,6 +2670,7 @@
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:one_pointer_one_module_scope:*'>
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:subcalls:*'>
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:member_accessors:*'>
+<meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:swizzles:*'>
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:same_pointer_read_and_write:*'>
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:aliasing_inside_function:*'>
 <meta name=variant content='?q=webgpu:shader,validation,functions,alias_analysis:two_atomic_pointers:*'>
@@ -2741,6 +2889,7 @@
 <meta name=variant content='?q=webgpu:shader,validation,statement,phony:rhs_constructible:*'>
 <meta name=variant content='?q=webgpu:shader,validation,statement,phony:rhs_with_decl:*'>
 <meta name=variant content='?q=webgpu:shader,validation,statement,phony:parse:*'>
+<meta name=variant content='?q=webgpu:shader,validation,statement,phony:module_scope:*'>
 <meta name=variant content='?q=webgpu:shader,validation,statement,return:return_missing_value:*'>
 <meta name=variant content='?q=webgpu:shader,validation,statement,return:return_unexpected_value:*'>
 <meta name=variant content='?q=webgpu:shader,validation,statement,return:return_type_match:*'>
@@ -2870,6 +3019,7 @@
 <meta name=variant content='?q=webgpu:web_platform,external_texture,video:importExternalTexture,sample_non_YUV_video_frame:*'>
 <meta name=variant content='?q=webgpu:web_platform,external_texture,video:importExternalTexture,sampleWithVideoFrameWithVisibleRectParam:*'>
 <meta name=variant content='?q=webgpu:web_platform,external_texture,video:importExternalTexture,compute:*'>
+<meta name=variant content='?q=webgpu:web_platform,external_texture,video:importExternalTexture,cameraCapture:*'>
 <meta name=variant content='?q=webgpu:web_platform,worker,worker:dedicated_worker:*'>
 <meta name=variant content='?q=webgpu:web_platform,worker,worker:shared_worker:*'>
 <meta name=variant content='?q=webgpu:web_platform,worker,worker:service_worker:*'>

tests/wpt/webgpu/tests/webgpu/resources/cache/hashes.json

@@ -1,112 +1,112 @@
 {
-  "webgpu/shader/execution/binary/af_addition.bin": "338b5b67",
-  "webgpu/shader/execution/binary/af_logical.bin": "3b2aceb8",
-  "webgpu/shader/execution/binary/af_division.bin": "a77dc4c0",
-  "webgpu/shader/execution/binary/af_matrix_addition.bin": "136a7fbb",
-  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "90f2c731",
-  "webgpu/shader/execution/binary/af_multiplication.bin": "35ba40b9",
-  "webgpu/shader/execution/binary/af_remainder.bin": "41582f85",
-  "webgpu/shader/execution/binary/af_subtraction.bin": "a41420b2",
-  "webgpu/shader/execution/binary/f16_addition.bin": "ef10ca66",
-  "webgpu/shader/execution/binary/f16_logical.bin": "4bf24ca5",
-  "webgpu/shader/execution/binary/f16_division.bin": "f826b6ba",
-  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "a910ddb0",
-  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "9458671c",
-  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "36be05d3",
-  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "8aa6a88a",
-  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "38282a11",
-  "webgpu/shader/execution/binary/f16_multiplication.bin": "62f91819",
-  "webgpu/shader/execution/binary/f16_remainder.bin": "f829bb65",
-  "webgpu/shader/execution/binary/f16_subtraction.bin": "82d4e231",
-  "webgpu/shader/execution/binary/f32_addition.bin": "9b0a0c50",
-  "webgpu/shader/execution/binary/f32_logical.bin": "b75af25a",
-  "webgpu/shader/execution/binary/f32_division.bin": "f6d7832f",
-  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "3317c75b",
-  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "c6f990c8",
-  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "b091a702",
-  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "2d12a16b",
-  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "e1217524",
-  "webgpu/shader/execution/binary/f32_multiplication.bin": "19774fb3",
-  "webgpu/shader/execution/binary/f32_remainder.bin": "fd94bb9a",
-  "webgpu/shader/execution/binary/f32_subtraction.bin": "dba7cd7a",
-  "webgpu/shader/execution/binary/i32_arithmetic.bin": "e3b317e1",
-  "webgpu/shader/execution/binary/i32_comparison.bin": "63fa9be8",
-  "webgpu/shader/execution/binary/u32_arithmetic.bin": "e8b4008c",
-  "webgpu/shader/execution/binary/u32_comparison.bin": "d472fd61",
-  "webgpu/shader/execution/abs.bin": "631d932d",
-  "webgpu/shader/execution/acos.bin": "afcafcb1",
-  "webgpu/shader/execution/acosh.bin": "4b30eb95",
-  "webgpu/shader/execution/asin.bin": "c850c13d",
-  "webgpu/shader/execution/asinh.bin": "66a6acc0",
-  "webgpu/shader/execution/atan.bin": "2aabbb53",
-  "webgpu/shader/execution/atan2.bin": "82dd926a",
-  "webgpu/shader/execution/atanh.bin": "b98c937c",
-  "webgpu/shader/execution/bitcast.bin": "5daaee1b",
-  "webgpu/shader/execution/ceil.bin": "d0c32cf4",
-  "webgpu/shader/execution/clamp.bin": "4d1fc26a",
-  "webgpu/shader/execution/cos.bin": "dc837ae2",
-  "webgpu/shader/execution/cosh.bin": "d9e90580",
-  "webgpu/shader/execution/cross.bin": "ce7979f",
-  "webgpu/shader/execution/degrees.bin": "1436a196",
-  "webgpu/shader/execution/determinant.bin": "f36f1fa1",
-  "webgpu/shader/execution/distance.bin": "5103f8bd",
-  "webgpu/shader/execution/dot.bin": "4514172c",
-  "webgpu/shader/execution/exp.bin": "f41150bd",
-  "webgpu/shader/execution/exp2.bin": "19c494e",
-  "webgpu/shader/execution/faceForward.bin": "27b6e4a7",
-  "webgpu/shader/execution/floor.bin": "5bb5098b",
-  "webgpu/shader/execution/fma.bin": "daace9a4",
-  "webgpu/shader/execution/fract.bin": "be5f0334",
-  "webgpu/shader/execution/frexp.bin": "c9efaf7c",
-  "webgpu/shader/execution/inverseSqrt.bin": "8a50b907",
-  "webgpu/shader/execution/ldexp.bin": "cb4cea21",
-  "webgpu/shader/execution/length.bin": "a1b9fbeb",
-  "webgpu/shader/execution/log.bin": "9f2eb7c3",
-  "webgpu/shader/execution/log2.bin": "9ee7d861",
-  "webgpu/shader/execution/max.bin": "11e4608e",
-  "webgpu/shader/execution/min.bin": "7a084c44",
-  "webgpu/shader/execution/mix.bin": "7b892a4f",
-  "webgpu/shader/execution/modf.bin": "b3bf26d7",
-  "webgpu/shader/execution/normalize.bin": "18eba01d",
-  "webgpu/shader/execution/pack2x16float.bin": "82df446e",
-  "webgpu/shader/execution/pow.bin": "d3a05344",
-  "webgpu/shader/execution/quantizeToF16.bin": "7793770e",
-  "webgpu/shader/execution/radians.bin": "582c1f6b",
-  "webgpu/shader/execution/reflect.bin": "9161d6e5",
-  "webgpu/shader/execution/refract.bin": "817b59aa",
-  "webgpu/shader/execution/round.bin": "cb881aa2",
-  "webgpu/shader/execution/saturate.bin": "3716605e",
-  "webgpu/shader/execution/sign.bin": "549ac92f",
-  "webgpu/shader/execution/sin.bin": "5ec5bcb7",
-  "webgpu/shader/execution/sinh.bin": "62f6b736",
-  "webgpu/shader/execution/smoothstep.bin": "aa97768",
-  "webgpu/shader/execution/sqrt.bin": "d0a134ce",
-  "webgpu/shader/execution/step.bin": "b8035bb9",
-  "webgpu/shader/execution/tan.bin": "b34366cd",
-  "webgpu/shader/execution/tanh.bin": "8f5edddc",
-  "webgpu/shader/execution/transpose.bin": "1aa2de65",
-  "webgpu/shader/execution/trunc.bin": "cf43e3f7",
-  "webgpu/shader/execution/unpack2x16float.bin": "57ea7c02",
-  "webgpu/shader/execution/unpack2x16snorm.bin": "17fd3f86",
-  "webgpu/shader/execution/unpack2x16unorm.bin": "fc68bc4b",
-  "webgpu/shader/execution/unpack4x8snorm.bin": "fef504c1",
-  "webgpu/shader/execution/unpack4x8unorm.bin": "e8d8de93",
-  "webgpu/shader/execution/unary/af_arithmetic.bin": "14c0612a",
-  "webgpu/shader/execution/unary/af_assignment.bin": "3ad4afc",
-  "webgpu/shader/execution/unary/bool_conversion.bin": "15f7f3fb",
-  "webgpu/shader/execution/unary/f16_arithmetic.bin": "4a20db6d",
-  "webgpu/shader/execution/unary/f16_conversion.bin": "31f72f5a",
-  "webgpu/shader/execution/unary/f32_arithmetic.bin": "f1c311cb",
-  "webgpu/shader/execution/unary/f32_conversion.bin": "7539cdb3",
-  "webgpu/shader/execution/unary/i32_arithmetic.bin": "de945eec",
-  "webgpu/shader/execution/unary/i32_conversion.bin": "1728a03e",
-  "webgpu/shader/execution/unary/u32_conversion.bin": "9e6ca0ce",
-  "webgpu/shader/execution/unary/ai_assignment.bin": "1fd685a2",
-  "webgpu/shader/execution/binary/ai_arithmetic.bin": "90e651f4",
-  "webgpu/shader/execution/unary/ai_arithmetic.bin": "ba31d178",
-  "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "bc8b52ef",
-  "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "54edf6a2",
-  "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "43b036b1",
-  "webgpu/shader/execution/derivatives.bin": "65c15fc3",
-  "webgpu/shader/execution/fwidth.bin": "cc91c875"
+  "webgpu/shader/execution/binary/af_addition.bin": "d0c1b760",
+  "webgpu/shader/execution/binary/af_logical.bin": "ca60ce77",
+  "webgpu/shader/execution/binary/af_division.bin": "47ae1ca1",
+  "webgpu/shader/execution/binary/af_matrix_addition.bin": "afaf9bae",
+  "webgpu/shader/execution/binary/af_matrix_subtraction.bin": "42433bf3",
+  "webgpu/shader/execution/binary/af_multiplication.bin": "babfc501",
+  "webgpu/shader/execution/binary/af_remainder.bin": "19995293",
+  "webgpu/shader/execution/binary/af_subtraction.bin": "62f090b9",
+  "webgpu/shader/execution/binary/f16_addition.bin": "540ae334",
+  "webgpu/shader/execution/binary/f16_logical.bin": "c1f09c30",
+  "webgpu/shader/execution/binary/f16_division.bin": "b4eabc05",
+  "webgpu/shader/execution/binary/f16_matrix_addition.bin": "6b9113b",
+  "webgpu/shader/execution/binary/f16_matrix_matrix_multiplication.bin": "a7362ff1",
+  "webgpu/shader/execution/binary/f16_matrix_scalar_multiplication.bin": "4ac4e5bb",
+  "webgpu/shader/execution/binary/f16_matrix_subtraction.bin": "93d4d43a",
+  "webgpu/shader/execution/binary/f16_matrix_vector_multiplication.bin": "beed89d5",
+  "webgpu/shader/execution/binary/f16_multiplication.bin": "6b5f0d51",
+  "webgpu/shader/execution/binary/f16_remainder.bin": "a1f499b3",
+  "webgpu/shader/execution/binary/f16_subtraction.bin": "61a571d5",
+  "webgpu/shader/execution/binary/f32_addition.bin": "fa6cc596",
+  "webgpu/shader/execution/binary/f32_logical.bin": "2b155b60",
+  "webgpu/shader/execution/binary/f32_division.bin": "243c9ce6",
+  "webgpu/shader/execution/binary/f32_matrix_addition.bin": "d3bc6ed6",
+  "webgpu/shader/execution/binary/f32_matrix_matrix_multiplication.bin": "2a4c1527",
+  "webgpu/shader/execution/binary/f32_matrix_scalar_multiplication.bin": "d695442",
+  "webgpu/shader/execution/binary/f32_matrix_subtraction.bin": "b306b19",
+  "webgpu/shader/execution/binary/f32_matrix_vector_multiplication.bin": "aac6cbfd",
+  "webgpu/shader/execution/binary/f32_multiplication.bin": "a21303f5",
+  "webgpu/shader/execution/binary/f32_remainder.bin": "79e462a1",
+  "webgpu/shader/execution/binary/f32_subtraction.bin": "4e6bbf38",
+  "webgpu/shader/execution/binary/i32_arithmetic.bin": "167760cc",
+  "webgpu/shader/execution/binary/i32_comparison.bin": "6a9f856a",
+  "webgpu/shader/execution/binary/u32_arithmetic.bin": "ac424b44",
+  "webgpu/shader/execution/binary/u32_comparison.bin": "a9e71302",
+  "webgpu/shader/execution/abs.bin": "a42729c4",
+  "webgpu/shader/execution/acos.bin": "664a5662",
+  "webgpu/shader/execution/acosh.bin": "d3fb8eb0",
+  "webgpu/shader/execution/asin.bin": "5a4f5b9e",
+  "webgpu/shader/execution/asinh.bin": "3ce3fe4d",
+  "webgpu/shader/execution/atan.bin": "759d432",
+  "webgpu/shader/execution/atan2.bin": "95008607",
+  "webgpu/shader/execution/atanh.bin": "569bd1b6",
+  "webgpu/shader/execution/bitcast.bin": "4329e501",
+  "webgpu/shader/execution/ceil.bin": "55cc76e5",
+  "webgpu/shader/execution/clamp.bin": "d580a273",
+  "webgpu/shader/execution/cos.bin": "3107bc4b",
+  "webgpu/shader/execution/cosh.bin": "d36c86cc",
+  "webgpu/shader/execution/cross.bin": "e48c39ba",
+  "webgpu/shader/execution/degrees.bin": "f74b63d2",
+  "webgpu/shader/execution/determinant.bin": "f07e1160",
+  "webgpu/shader/execution/distance.bin": "93156a89",
+  "webgpu/shader/execution/dot.bin": "4e2fe407",
+  "webgpu/shader/execution/exp.bin": "3b269b18",
+  "webgpu/shader/execution/exp2.bin": "7aeeeaf6",
+  "webgpu/shader/execution/faceForward.bin": "451ffbd8",
+  "webgpu/shader/execution/floor.bin": "37131d74",
+  "webgpu/shader/execution/fma.bin": "30111350",
+  "webgpu/shader/execution/fract.bin": "5ef13392",
+  "webgpu/shader/execution/frexp.bin": "da764bc0",
+  "webgpu/shader/execution/inverseSqrt.bin": "6ff34703",
+  "webgpu/shader/execution/ldexp.bin": "5016cec9",
+  "webgpu/shader/execution/length.bin": "f236d2e7",
+  "webgpu/shader/execution/log.bin": "1c54f128",
+  "webgpu/shader/execution/log2.bin": "e44e2370",
+  "webgpu/shader/execution/max.bin": "eb4c1901",
+  "webgpu/shader/execution/min.bin": "f8c70a2b",
+  "webgpu/shader/execution/mix.bin": "df3b3f62",
+  "webgpu/shader/execution/modf.bin": "b600b26f",
+  "webgpu/shader/execution/normalize.bin": "7af3a3d2",
+  "webgpu/shader/execution/pack2x16float.bin": "7c67b10e",
+  "webgpu/shader/execution/pow.bin": "ee37f4ba",
+  "webgpu/shader/execution/quantizeToF16.bin": "a7a65754",
+  "webgpu/shader/execution/radians.bin": "51d423b9",
+  "webgpu/shader/execution/reflect.bin": "3ba4eda6",
+  "webgpu/shader/execution/refract.bin": "13fc4914",
+  "webgpu/shader/execution/round.bin": "9155b88b",
+  "webgpu/shader/execution/saturate.bin": "73cecf71",
+  "webgpu/shader/execution/sign.bin": "68d61a83",
+  "webgpu/shader/execution/sin.bin": "44219876",
+  "webgpu/shader/execution/sinh.bin": "158d261d",
+  "webgpu/shader/execution/smoothstep.bin": "7129c56b",
+  "webgpu/shader/execution/sqrt.bin": "9aaaf8aa",
+  "webgpu/shader/execution/step.bin": "85858027",
+  "webgpu/shader/execution/tan.bin": "dbbda634",
+  "webgpu/shader/execution/tanh.bin": "8c540d5c",
+  "webgpu/shader/execution/transpose.bin": "a676fc9a",
+  "webgpu/shader/execution/trunc.bin": "35ab398d",
+  "webgpu/shader/execution/unpack2x16float.bin": "eb9294c9",
+  "webgpu/shader/execution/unpack2x16snorm.bin": "7208eb73",
+  "webgpu/shader/execution/unpack2x16unorm.bin": "20d9669b",
+  "webgpu/shader/execution/unpack4x8snorm.bin": "c77e1a72",
+  "webgpu/shader/execution/unpack4x8unorm.bin": "d80caf66",
+  "webgpu/shader/execution/unary/af_arithmetic.bin": "963c3185",
+  "webgpu/shader/execution/unary/af_assignment.bin": "9e8a3b3f",
+  "webgpu/shader/execution/unary/bool_conversion.bin": "eee7a40c",
+  "webgpu/shader/execution/unary/f16_arithmetic.bin": "aaea9f75",
+  "webgpu/shader/execution/unary/f16_conversion.bin": "5b26998a",
+  "webgpu/shader/execution/unary/f32_arithmetic.bin": "65dfc2ac",
+  "webgpu/shader/execution/unary/f32_conversion.bin": "cd874be3",
+  "webgpu/shader/execution/unary/i32_arithmetic.bin": "af4c0e43",
+  "webgpu/shader/execution/unary/i32_conversion.bin": "5b6e4d9",
+  "webgpu/shader/execution/unary/u32_conversion.bin": "229649a6",
+  "webgpu/shader/execution/unary/ai_assignment.bin": "8efcf261",
+  "webgpu/shader/execution/binary/ai_arithmetic.bin": "a57ee284",
+  "webgpu/shader/execution/unary/ai_arithmetic.bin": "948016b6",
+  "webgpu/shader/execution/binary/af_matrix_matrix_multiplication.bin": "52c24212",
+  "webgpu/shader/execution/binary/af_matrix_scalar_multiplication.bin": "256556e1",
+  "webgpu/shader/execution/binary/af_matrix_vector_multiplication.bin": "38085521",
+  "webgpu/shader/execution/derivatives.bin": "f38a38ff",
+  "webgpu/shader/execution/fwidth.bin": "4e9fc55d"
 }

tests/wpt/webgpu/tests/webgpu/webgpu/api/operation/rendering/color_target_state.spec.js

@@ -9,7 +9,11 @@ TODO:
 - ?
 `;import { makeTestGroup } from '../../../../common/framework/test_group.js';
 import { assert, unreachable } from '../../../../common/util/util.js';
-import { kBlendFactors, kBlendOperations } from '../../../capability_info.js';
+import {
+  IsDualSourceBlendingFactor,
+  kBlendFactors,
+  kBlendOperations } from
+'../../../capability_info.js';
 import { GPUConst } from '../../../constants.js';
 import { kRegularTextureFormats, kTextureFormatInfo } from '../../../format_info.js';
 import { GPUTest, TextureTestMixin } from '../../../gpu_test.js';
@@ -85,6 +89,7 @@ f)
 
 function computeBlendFactor(
 src,
+src1,
 dst,
 blendColor,
 factor)
@@ -120,6 +125,14 @@ factor)
     case 'one-minus-constant':
       assert(blendColor !== undefined);
       return mapColor(blendColor, (v) => 1 - v);
+    case 'src1':
+      return { ...src1 };
+    case 'one-minus-src1':
+      return mapColor(src1, (v) => 1 - v);
+    case 'src1-alpha':
+      return mapColor(src1, () => src1.a);
+    case 'one-minus-src1-alpha':
+      return mapColor(src1, () => 1 - src1.a);
     default:
       unreachable();
   }
@@ -174,6 +187,7 @@ filter((t) => {
   return true;
 }).
 combine('srcColor', [{ r: 0.11, g: 0.61, b: 0.81, a: 0.44 }]).
+combine('srcColor1', [{ r: 0.22, g: 0.41, b: 0.51, a: 0.33 }]).
 combine('dstColor', [
 { r: 0.51, g: 0.22, b: 0.71, a: 0.33 },
 { r: 0.09, g: 0.73, b: 0.93, a: 0.81 }]
@@ -187,14 +201,35 @@ expand('blendConstant', (p) => {
   return needsBlendConstant ? [{ r: 0.91, g: 0.82, b: 0.73, a: 0.64 }] : [undefined];
 })
 ).
+beforeAllSubcases((t) => {
+  if (
+  IsDualSourceBlendingFactor(t.params.srcFactor) ||
+  IsDualSourceBlendingFactor(t.params.dstFactor))
+  {
+    t.selectDeviceOrSkipTestCase('dual-source-blending');
+  }
+}).
 fn((t) => {
   const textureFormat = 'rgba16float';
   const srcColor = t.params.srcColor;
+  const srcColor1 = t.params.srcColor1;
   const dstColor = t.params.dstColor;
   const blendConstant = t.params.blendConstant;
 
-  const srcFactor = computeBlendFactor(srcColor, dstColor, blendConstant, t.params.srcFactor);
-  const dstFactor = computeBlendFactor(srcColor, dstColor, blendConstant, t.params.dstFactor);
+  const srcFactor = computeBlendFactor(
+    srcColor,
+    srcColor1,
+    dstColor,
+    blendConstant,
+    t.params.srcFactor
+  );
+  const dstFactor = computeBlendFactor(
+    srcColor,
+    srcColor1,
+    dstColor,
+    blendConstant,
+    t.params.dstFactor
+  );
 
   const expectedColor = computeBlendOperation(
     srcColor,
@@ -215,6 +250,10 @@ fn((t) => {
       break;
   }
 
+  const useBlendSrc1 =
+  IsDualSourceBlendingFactor(t.params.srcFactor) ||
+  IsDualSourceBlendingFactor(t.params.dstFactor);
+
   const pipeline = t.device.createRenderPipeline({
     layout: 'auto',
     fragment: {
@@ -236,13 +275,24 @@ fn((t) => {
 
       module: t.device.createShaderModule({
         code: `
+${useBlendSrc1 ? 'enable dual_source_blending;' : ''}
+
 struct Uniform {
-  color: vec4<f32>
+  color: vec4f,
+  blend: vec4f,
 };
 @group(0) @binding(0) var<uniform> u : Uniform;
 
-@fragment fn main() -> @location(0) vec4<f32> {
-  return u.color;
+struct FragOutput {
+  @location(0) ${useBlendSrc1 ? '@blend_src(0)' : ''} color : vec4f,
+  ${useBlendSrc1 ? '@location(0) @blend_src(1) blend : vec4f,' : ''}
+}
+
+@fragment fn main() ->FragOutput {
+  var fragOutput : FragOutput;
+  fragOutput.color = u.color;
+  ${useBlendSrc1 ? 'fragOutput.blend = u.blend;' : ''}
+  return fragOutput;
 }
           `
       }),
@@ -293,7 +343,16 @@ struct Uniform {
         binding: 0,
         resource: {
           buffer: t.makeBufferWithContents(
-            new Float32Array([srcColor.r, srcColor.g, srcColor.b, srcColor.a]),
+            new Float32Array([
+            srcColor.r,
+            srcColor.g,
+            srcColor.b,
+            srcColor.a,
+            srcColor1.r,
+            srcColor1.g,
+            srcColor1.b,
+            srcColor1.a]
+            ),
             GPUBufferUsage.UNIFORM
           )
         }

tests/wpt/webgpu/tests/webgpu/webgpu/api/operation/shader_module/compilation_info.spec.js

@@ -3,7 +3,6 @@
 **/export const description = `
 ShaderModule CompilationInfo tests.
 `;import { makeTestGroup } from '../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../common/util/data_tables.js';
 import { assert } from '../../../../common/util/util.js';
 import { GPUTest } from '../../../gpu_test.js';
 
@@ -79,66 +78,24 @@ const kInvalidShaderSources = [
 
 const kAllShaderSources = [...kValidShaderSources, ...kInvalidShaderSources];
 
-// This is the source the sourcemap refers to.
-const kOriginalSource = new Array(20).
-fill(0).
-map((_, i) => `original line ${i}`).
-join('\n');
-
-const kSourceMaps = {
-  none: undefined,
-  empty: {},
-  // A valid source map. It maps `unknown` on lines 4 and line 5 to
-  // `wasUnknown` from lines 20, 21 respectively
-  valid: {
-    version: 3,
-    sources: ['myCode'],
-    sourcesContent: [kOriginalSource],
-    names: ['myMain', 'wasUnknown'],
-    mappings: ';kBAYkCA,OACd;SAElB;gBAKOC;gBACAA'
-  },
-  // not a valid sourcemap
-  invalid: {
-    version: -123,
-    notAnything: {}
-  },
-  // The correct format but this data is for lines 11,12 even
-  // though the source only has 5 or 6 lines
-  nonMatching: {
-    version: 3,
-    sources: ['myCode'],
-    sourcesContent: [kOriginalSource],
-    names: ['myMain'],
-    mappings: ';;;;;;;;;;kBAYkCA,OACd;SAElB'
-  }
-};
-const kSourceMapsKeys = keysOf(kSourceMaps);
-
 g.test('getCompilationInfo_returns').
 desc(
   `
     Test that getCompilationInfo() can be called on any ShaderModule.
 
-    Note: sourcemaps are not used in the WebGPU API. We are only testing that
-    browser that happen to use them don't fail or crash if the sourcemap is
-    bad or invalid.
-
     - Test for both valid and invalid shader modules.
     - Test for shader modules containing only ASCII and those containing unicode characters.
     - Test that the compilation info for valid shader modules contains no errors.
     - Test that the compilation info for invalid shader modules contains at least one error.`
 ).
-params((u) =>
-u.combineWithParams(kAllShaderSources).beginSubcases().combine('sourceMapName', kSourceMapsKeys)
-).
+params((u) => u.combineWithParams(kAllShaderSources)).
 fn(async (t) => {
-  const { _code, valid, sourceMapName } = t.params;
+  const { _code, valid } = t.params;
 
   const shaderModule = t.expectGPUError(
     'validation',
     () => {
-      const sourceMap = kSourceMaps[sourceMapName];
-      return t.device.createShaderModule({ code: _code, ...(sourceMap && { sourceMap }) });
+      return t.device.createShaderModule({ code: _code });
     },
     !valid
   );
@@ -171,25 +128,15 @@ desc(
     Test that line numbers reported by compilationInfo either point at an appropriate line and
     position or at 0:0, indicating an unknown position.
 
-    Note: sourcemaps are not used in the WebGPU API. We are only testing that
-    browser that happen to use them don't fail or crash if the sourcemap is
-    bad or invalid.
-
     - Test for invalid shader modules containing containing at least one error.
     - Test for shader modules containing only ASCII and those containing unicode characters.`
 ).
-params((u) =>
-u.
-combineWithParams(kInvalidShaderSources).
-beginSubcases().
-combine('sourceMapName', kSourceMapsKeys)
-).
+params((u) => u.combineWithParams(kInvalidShaderSources)).
 fn(async (t) => {
-  const { _code, _errorLine, _errorLinePos, sourceMapName } = t.params;
+  const { _code, _errorLine, _errorLinePos } = t.params;
 
   const shaderModule = t.expectGPUError('validation', () => {
-    const sourceMap = kSourceMaps[sourceMapName];
-    return t.device.createShaderModule({ code: _code, ...(sourceMap && { sourceMap }) });
+    return t.device.createShaderModule({ code: _code });
   });
 
   const info = await shaderModule.getCompilationInfo();
@@ -232,24 +179,17 @@ g.test('offset_and_length').
 desc(
   `Test that message offsets and lengths are valid and align with any reported lineNum and linePos.
 
-     Note: sourcemaps are not used in the WebGPU API. We are only testing that
-     browser that happen to use them don't fail or crash if the sourcemap is
-     bad or invalid.
-
     - Test for valid and invalid shader modules.
     - Test for shader modules containing only ASCII and those containing unicode characters.`
 ).
-params((u) =>
-u.combineWithParams(kAllShaderSources).beginSubcases().combine('sourceMapName', kSourceMapsKeys)
-).
+params((u) => u.combineWithParams(kAllShaderSources)).
 fn(async (t) => {
-  const { _code, valid, sourceMapName } = t.params;
+  const { _code, valid } = t.params;
 
   const shaderModule = t.expectGPUError(
     'validation',
     () => {
-      const sourceMap = kSourceMaps[sourceMapName];
-      return t.device.createShaderModule({ code: _code, ...(sourceMap && { sourceMap }) });
+      return t.device.createShaderModule({ code: _code });
     },
     !valid
   );

tests/wpt/webgpu/tests/webgpu/webgpu/api/operation/texture_view/write.spec.js

@@ -36,6 +36,9 @@ const kTextureViewWriteMethods = [
 
 
 
+const kTextureViewUsageMethods = ['inherit', 'minimal'];
+
+
 // Src color values to read from a shader array.
 const kColorsFloat = [
 { R: 1.0, G: 0.0, B: 0.0, A: 0.8 },
@@ -271,6 +274,22 @@ sampleCount)
   return expectedTexelView;
 }
 
+function getTextureViewUsage(
+viewUsageMethod,
+minimalUsageForTest)
+{
+  switch (viewUsageMethod) {
+    case 'inherit':
+      return 0;
+
+    case 'minimal':
+      return minimalUsageForTest;
+
+    default:
+      unreachable();
+  }
+}
+
 g.test('format').
 desc(
   `Views of every allowed format.
@@ -280,6 +299,7 @@ Read values from color array in the shader, and write it to the texture view via
 - x= every texture format
 - x= sampleCount {1, 4} if valid
 - x= every possible view write method (see above)
+- x= inherited or minimal texture view usage
 
 TODO: Test sampleCount > 1 for 'render-pass-store' after extending copySinglePixelTextureToBufferUsingComputePass
       to read multiple pixels from multisampled textures. [1]
@@ -317,7 +337,8 @@ filter(({ format, method, sampleCount }) => {
       return !!info.colorRender?.resolve && sampleCount === 1;
   }
   return true;
-})
+}).
+combine('viewUsageMethod', kTextureViewUsageMethods)
 ).
 beforeAllSubcases((t) => {
   const { format, method } = t.params;
@@ -332,13 +353,12 @@ beforeAllSubcases((t) => {
   }
 }).
 fn((t) => {
-  const { format, method, sampleCount } = t.params;
+  const { format, method, sampleCount, viewUsageMethod } = t.params;
 
-  const usage =
-  GPUTextureUsage.COPY_SRC | (
-  method.includes('storage') ?
+  const textureUsageForMethod = method.includes('storage') ?
   GPUTextureUsage.STORAGE_BINDING :
-  GPUTextureUsage.RENDER_ATTACHMENT);
+  GPUTextureUsage.RENDER_ATTACHMENT;
+  const usage = GPUTextureUsage.COPY_SRC | textureUsageForMethod;
 
   const texture = t.createTextureTracked({
     format,
@@ -347,7 +367,9 @@ fn((t) => {
     sampleCount
   });
 
-  const view = texture.createView();
+  const view = texture.createView({
+    usage: getTextureViewUsage(viewUsageMethod, textureUsageForMethod)
+  });
   const expectedTexelView = writeTextureAndGetExpectedTexelView(
     t,
     method,

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/features/clip_distances.spec.js

@@ -0,0 +1,161 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/import { range } from '../../../../../common/util/util.js';import { align } from '../../../../util/math.js';import { kMaximumLimitBaseParams, makeLimitTestGroup } from '../limits/limit_utils.js';
+
+function getPipelineDescriptorWithClipDistances(
+device,
+interStageShaderVariables,
+pointList,
+clipDistances,
+startLocation = 0)
+{
+  const vertexOutputVariables =
+  interStageShaderVariables - (pointList ? 1 : 0) - align(clipDistances, 4) / 4;
+  const maxVertexOutputVariables =
+  device.limits.maxInterStageShaderVariables - (pointList ? 1 : 0) - align(clipDistances, 4) / 4;
+
+  const varyings = `
+      ${range(
+    vertexOutputVariables,
+    (i) => `@location(${i + startLocation}) v4_${i + startLocation}: vec4f,`
+  ).join('\n')}
+  `;
+
+  const code = `
+    // test value                        : ${interStageShaderVariables}
+    // maxInterStageShaderVariables     : ${device.limits.maxInterStageShaderVariables}
+    // num variables in vertex shader : ${vertexOutputVariables}${
+  pointList ? ' + point-list' : ''
+  }${
+  clipDistances > 0 ?
+  ` + ${align(clipDistances, 4) / 4} (clip_distances[${clipDistances}])` :
+  ''
+  }
+    // maxInterStageVariables:           : ${maxVertexOutputVariables}
+    // num used inter stage variables    : ${vertexOutputVariables}
+    // vertex output start location      : ${startLocation}
+
+    enable clip_distances;
+
+    struct VSOut {
+      @builtin(position) p: vec4f,
+      ${varyings}
+      ${
+  clipDistances > 0 ?
+  `@builtin(clip_distances) clipDistances: array<f32, ${clipDistances}>,` :
+  ''
+  }
+    }
+    struct FSIn {
+      ${varyings}
+    }
+    struct FSOut {
+      @location(0) color: vec4f,
+    }
+    @vertex fn vs() -> VSOut {
+      var o: VSOut;
+      o.p = vec4f(0);
+      return o;
+    }
+    @fragment fn fs(i: FSIn) -> FSOut {
+      var o: FSOut;
+      o.color = vec4f(0);
+      return o;
+    }
+  `;
+  const module = device.createShaderModule({ code });
+  const pipelineDescriptor = {
+    layout: 'auto',
+    primitive: {
+      topology: pointList ? 'point-list' : 'triangle-list'
+    },
+    vertex: {
+      module
+    },
+    fragment: {
+      module,
+      targets: [
+      {
+        format: 'rgba8unorm'
+      }]
+
+    }
+  };
+  return pipelineDescriptor;
+}
+
+const limit = 'maxInterStageShaderVariables';
+export const { g, description } = makeLimitTestGroup(limit);
+
+g.test('createRenderPipeline,at_over').
+desc(`Test using at and over ${limit} limit with clip_distances in createRenderPipeline(Async)`).
+params(
+  kMaximumLimitBaseParams.
+  combine('async', [false, true]).
+  combine('pointList', [false, true]).
+  combine('clipDistances', [1, 2, 3, 4, 5, 6, 7, 8])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('clip-distances');
+}).
+fn(async (t) => {
+  const { limitTest, testValueName, async, pointList, clipDistances } = t.params;
+  await t.testDeviceWithRequestedMaximumLimits(
+    limitTest,
+    testValueName,
+    async ({ device, testValue, shouldError }) => {
+      const pipelineDescriptor = getPipelineDescriptorWithClipDistances(
+        device,
+        testValue,
+        pointList,
+        clipDistances
+      );
+
+      await t.testCreateRenderPipeline(pipelineDescriptor, async, shouldError);
+    },
+    undefined,
+    ['clip-distances']
+  );
+});
+
+g.test('createRenderPipeline,max_vertex_output_location').
+desc(`Test using clip_distances will limit the maximum value of vertex output location`).
+params((u) =>
+u.
+combine('pointList', [false, true]).
+combine('clipDistances', [1, 2, 3, 4, 5, 6, 7, 8]).
+combine('startLocation', [0, 1, 2])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('clip-distances');
+}).
+fn(async (t) => {
+  const { pointList, clipDistances, startLocation } = t.params;
+
+  const maxInterStageShaderVariables = t.adapter.limits.maxInterStageShaderVariables;
+  const deviceInTest = await t.requestDeviceTracked(t.adapter, {
+    requiredFeatures: ['clip-distances'],
+    requiredLimits: {
+      maxInterStageShaderVariables: t.adapter.limits.maxInterStageShaderVariables
+    }
+  });
+  const pipelineDescriptor = getPipelineDescriptorWithClipDistances(
+    deviceInTest,
+    maxInterStageShaderVariables,
+    pointList,
+    clipDistances,
+    startLocation
+  );
+  const vertexOutputVariables =
+  maxInterStageShaderVariables - (pointList ? 1 : 0) - align(clipDistances, 4) / 4;
+  const maxLocationInTest = startLocation + vertexOutputVariables - 1;
+  const maxAllowedLocation = maxInterStageShaderVariables - 1 - align(clipDistances, 4) / 4;
+  const shouldError = maxLocationInTest > maxAllowedLocation;
+
+  deviceInTest.pushErrorScope('validation');
+  deviceInTest.createRenderPipeline(pipelineDescriptor);
+  const error = await deviceInTest.popErrorScope();
+  t.expect(!!error === shouldError, `${error?.message || 'no error when one was expected'}`);
+
+  deviceInTest.destroy();
+});

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/features/texture_formats.spec.js

@@ -5,6 +5,7 @@ Tests for capability checking for features enabling optional texture formats.
 `;import { makeTestGroup } from '../../../../../common/framework/test_group.js';
 import { getGPU } from '../../../../../common/util/navigator_gpu.js';
 import { assert } from '../../../../../common/util/util.js';
+import { kCanvasTextureFormats } from '../../../../capability_info.js';
 import { kAllTextureFormats, kTextureFormatInfo } from '../../../../format_info.js';
 import { kAllCanvasTypes, createCanvas } from '../../../../util/create_elements.js';
 import { ValidationTest } from '../../validation_test.js';
@@ -161,15 +162,15 @@ fn((t) => {
     usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST
   };
 
-  if (enable_required_feature) {
-    t.expectValidationError(() => {
-      ctx.configure(canvasConf);
-    });
-  } else {
-    t.shouldThrow('TypeError', () => {
-      ctx.configure(canvasConf);
-    });
-  }
+  const expectedError =
+  enable_required_feature &&
+  kCanvasTextureFormats.includes(format) ?
+  false :
+  'TypeError';
+
+  t.shouldThrow(expectedError, () => {
+    ctx.configure(canvasConf);
+  });
 });
 
 g.test('canvas_configuration_view_formats').

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/limit_utils.js

@@ -535,11 +535,16 @@ export class LimitTestsImpl extends GPUTestBase {
   limitTest,
   testValueName,
   fn,
-  extraLimits)
+  extraLimits,
+  extraFeatures = [])
   {
     assert(!this._device);
 
-    const deviceAndLimits = await this._getDeviceWithRequestedMaximumLimit(limitTest, extraLimits);
+    const deviceAndLimits = await this._getDeviceWithRequestedMaximumLimit(
+      limitTest,
+      extraLimits,
+      extraFeatures
+    );
     // If we request over the limit requestDevice will throw
     if (!deviceAndLimits) {
       return;

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/maxInterStageShaderComponents.spec.js

@@ -1,153 +0,0 @@
-/**
-* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
-**/import { range } from '../../../../../common/util/util.js';import { kMaximumLimitBaseParams, makeLimitTestGroup } from './limit_utils.js';
-
-function getPipelineDescriptor(
-device,
-testValue,
-pointList,
-frontFacing,
-sampleIndex,
-sampleMaskIn,
-sampleMaskOut)
-{
-  const success = testValue <= device.limits.maxInterStageShaderComponents;
-
-  const maxVertexOutputComponents =
-  device.limits.maxInterStageShaderComponents - (pointList ? 1 : 0);
-  const maxFragmentInputComponents =
-  device.limits.maxInterStageShaderComponents - (
-  frontFacing ? 1 : 0) - (
-  sampleIndex ? 1 : 0) - (
-  sampleMaskIn ? 1 : 0);
-  const maxOutputComponents = Math.min(maxVertexOutputComponents, maxFragmentInputComponents);
-  const maxInterStageVariables = Math.floor(maxOutputComponents / 4);
-  const maxUserDefinedVertexComponents = Math.floor(maxVertexOutputComponents / 4) * 4;
-  const maxUserDefinedFragmentComponents = Math.floor(maxFragmentInputComponents / 4) * 4;
-
-  const numInterStageVariables = success ? maxInterStageVariables : maxInterStageVariables + 1;
-  const numUserDefinedComponents = numInterStageVariables * 4;
-
-  const varyings = `
-      ${range(numInterStageVariables, (i) => `@location(${i}) v4_${i}: vec4f,`).join('\n')}
-  `;
-
-  const code = `
-    // test value                        : ${testValue}
-    // maxInterStageShaderComponents     : ${device.limits.maxInterStageShaderComponents}
-    // num components in vertex shader   : ${numUserDefinedComponents}${
-  pointList ? ' + point-list' : ''
-  }
-    // num components in fragment shader : ${numUserDefinedComponents}${
-  frontFacing ? ' + front-facing' : ''
-  }${sampleIndex ? ' + sample_index' : ''}${sampleMaskIn ? ' + sample_mask' : ''}
-    // maxUserDefinedVertexShaderOutputComponents   : ${maxUserDefinedVertexComponents}
-    // maxUserDefinedFragmentShaderInputComponents  : ${maxUserDefinedFragmentComponents}
-    // maxInterStageVariables:           : ${maxInterStageVariables}
-    // num used inter stage variables    : ${numInterStageVariables}
-
-    struct VSOut {
-      @builtin(position) p: vec4f,
-      ${varyings}
-    }
-    struct FSIn {
-      ${frontFacing ? '@builtin(front_facing) frontFacing: bool,' : ''}
-      ${sampleIndex ? '@builtin(sample_index) sampleIndex: u32,' : ''}
-      ${sampleMaskIn ? '@builtin(sample_mask) sampleMask: u32,' : ''}
-      ${varyings}
-    }
-    struct FSOut {
-      @location(0) color: vec4f,
-      ${sampleMaskOut ? '@builtin(sample_mask) sampleMask: u32,' : ''}
-    }
-    @vertex fn vs() -> VSOut {
-      var o: VSOut;
-      o.p = vec4f(0);
-      return o;
-    }
-    @fragment fn fs(i: FSIn) -> FSOut {
-      var o: FSOut;
-      o.color = vec4f(0);
-      return o;
-    }
-  `;
-  const module = device.createShaderModule({ code });
-  const pipelineDescriptor = {
-    layout: 'auto',
-    primitive: {
-      topology: pointList ? 'point-list' : 'triangle-list'
-    },
-    vertex: {
-      module,
-      entryPoint: 'vs'
-    },
-    fragment: {
-      module,
-      entryPoint: 'fs',
-      targets: [
-      {
-        format: 'rgba8unorm'
-      }]
-
-    }
-  };
-  return { pipelineDescriptor, code };
-}
-
-const limit = 'maxInterStageShaderComponents';
-export const { g, description } = makeLimitTestGroup(limit);
-
-g.test('createRenderPipeline,at_over').
-desc(`Test using at and over ${limit} limit in createRenderPipeline(Async)`).
-params(
-  kMaximumLimitBaseParams.
-  combine('async', [false, true]).
-  combine('pointList', [false, true]).
-  combine('frontFacing', [false, true]).
-  combine('sampleIndex', [false, true]).
-  combine('sampleMaskIn', [false, true]).
-  combine('sampleMaskOut', [false, true])
-).
-beforeAllSubcases((t) => {
-  if (t.isCompatibility) {
-    t.skipIf(
-      t.params.sampleMaskIn || t.params.sampleMaskOut,
-      'sample_mask not supported in compatibility mode'
-    );
-    t.skipIf(t.params.sampleIndex, 'sample_index not supported in compatibility mode');
-  }
-}).
-fn(async (t) => {
-  const {
-    limitTest,
-    testValueName,
-    async,
-    pointList,
-    frontFacing,
-    sampleIndex,
-    sampleMaskIn,
-    sampleMaskOut
-  } = t.params;
-  // Request the largest value of maxInterStageShaderVariables to allow the test using as many
-  // inter-stage shader components as possible without being limited by
-  // maxInterStageShaderVariables.
-  const extraLimits = { maxInterStageShaderVariables: 'adapterLimit' };
-  await t.testDeviceWithRequestedMaximumLimits(
-    limitTest,
-    testValueName,
-    async ({ device, testValue, shouldError }) => {
-      const { pipelineDescriptor, code } = getPipelineDescriptor(
-        device,
-        testValue,
-        pointList,
-        frontFacing,
-        sampleIndex,
-        sampleMaskIn,
-        sampleMaskOut
-      );
-
-      await t.testCreateRenderPipeline(pipelineDescriptor, async, shouldError, code);
-    },
-    extraLimits
-  );
-});

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/maxInterStageShaderVariables.spec.js

@@ -1,26 +1,86 @@
 /**
 * AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
-**/import { kMaximumLimitBaseParams, makeLimitTestGroup } from './limit_utils.js';function getPipelineDescriptor(device, testValue) {
+**/import { range } from '../../../../../common/util/util.js';import { kMaximumLimitBaseParams, makeLimitTestGroup } from './limit_utils.js';
+
+function getPipelineDescriptor(
+device,
+testValue,
+pointList,
+frontFacing,
+sampleIndex,
+sampleMaskIn,
+sampleMaskOut)
+{
+  const vertexOutputVariables = testValue - (pointList ? 1 : 0);
+  const fragmentInputVariables = testValue - (frontFacing || sampleIndex || sampleMaskIn ? 1 : 0);
+  const numInterStageVariables = Math.min(vertexOutputVariables, fragmentInputVariables);
+
+  const maxVertexOutputVariables = device.limits.maxInterStageShaderVariables - (pointList ? 1 : 0);
+  const maxFragmentInputVariables =
+  device.limits.maxInterStageShaderVariables - (
+  frontFacing || sampleIndex || sampleMaskIn ? 1 : 0);
+  const maxInterStageVariables = Math.min(maxVertexOutputVariables, maxFragmentInputVariables);
+
+  const varyings = `
+      ${range(numInterStageVariables, (i) => `@location(${i}) v4_${i}: vec4f,`).join('\n')}
+  `;
+
   const code = `
+    // test value                        : ${testValue}
+    // maxInterStageShaderVariables     : ${device.limits.maxInterStageShaderVariables}
+    // num variables in vertex shader : ${vertexOutputVariables}${pointList ? ' + point-list' : ''}
+    // num variables in fragment shader : ${fragmentInputVariables}${
+  frontFacing ? ' + front-facing' : ''
+  }${sampleIndex ? ' + sample_index' : ''}${sampleMaskIn ? ' + sample_mask' : ''}
+    // maxInterStageVariables:           : ${maxInterStageVariables}
+    // num used inter stage variables    : ${numInterStageVariables}
+
     struct VSOut {
       @builtin(position) p: vec4f,
-      @location(${testValue}) v: f32,
+      ${varyings}
+    }
+    struct FSIn {
+      ${frontFacing ? '@builtin(front_facing) frontFacing: bool,' : ''}
+      ${sampleIndex ? '@builtin(sample_index) sampleIndex: u32,' : ''}
+      ${sampleMaskIn ? '@builtin(sample_mask) sampleMask: u32,' : ''}
+      ${varyings}
+    }
+    struct FSOut {
+      @location(0) color: vec4f,
+      ${sampleMaskOut ? '@builtin(sample_mask) sampleMask: u32,' : ''}
     }
     @vertex fn vs() -> VSOut {
       var o: VSOut;
       o.p = vec4f(0);
-      o.v = 1.0;
+      return o;
+    }
+    @fragment fn fs(i: FSIn) -> FSOut {
+      var o: FSOut;
+      o.color = vec4f(0);
       return o;
     }
   `;
   const module = device.createShaderModule({ code });
-  return {
+  const pipelineDescriptor = {
     layout: 'auto',
+    primitive: {
+      topology: pointList ? 'point-list' : 'triangle-list'
+    },
     vertex: {
       module,
       entryPoint: 'vs'
+    },
+    fragment: {
+      module,
+      entryPoint: 'fs',
+      targets: [
+      {
+        format: 'rgba8unorm'
+      }]
+
     }
   };
+  return pipelineDescriptor;
 }
 
 const limit = 'maxInterStageShaderVariables';
@@ -28,15 +88,48 @@ export const { g, description } = makeLimitTestGroup(limit);
 
 g.test('createRenderPipeline,at_over').
 desc(`Test using at and over ${limit} limit in createRenderPipeline(Async)`).
-params(kMaximumLimitBaseParams.combine('async', [false, true])).
+params(
+  kMaximumLimitBaseParams.
+  combine('async', [false, true]).
+  combine('pointList', [false, true]).
+  combine('frontFacing', [false, true]).
+  combine('sampleIndex', [false, true]).
+  combine('sampleMaskIn', [false, true]).
+  combine('sampleMaskOut', [false, true])
+).
+beforeAllSubcases((t) => {
+  if (t.isCompatibility) {
+    t.skipIf(
+      t.params.sampleMaskIn || t.params.sampleMaskOut,
+      'sample_mask not supported in compatibility mode'
+    );
+    t.skipIf(t.params.sampleIndex, 'sample_index not supported in compatibility mode');
+  }
+}).
 fn(async (t) => {
-  const { limitTest, testValueName, async } = t.params;
+  const {
+    limitTest,
+    testValueName,
+    async,
+    pointList,
+    frontFacing,
+    sampleIndex,
+    sampleMaskIn,
+    sampleMaskOut
+  } = t.params;
   await t.testDeviceWithRequestedMaximumLimits(
     limitTest,
     testValueName,
     async ({ device, testValue, shouldError }) => {
-      const lastIndex = testValue - 1;
-      const pipelineDescriptor = getPipelineDescriptor(device, lastIndex);
+      const pipelineDescriptor = getPipelineDescriptor(
+        device,
+        testValue,
+        pointList,
+        frontFacing,
+        sampleIndex,
+        sampleMaskIn,
+        sampleMaskOut
+      );
 
       await t.testCreateRenderPipeline(pipelineDescriptor, async, shouldError);
     }

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/maxVertexAttributes.spec.js

@@ -18,7 +18,8 @@
         attributes: [{ shaderLocation: lastIndex, offset: 0, format: 'float32x4' }]
       }]
 
-    }
+    },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
   };
 }
 

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/maxVertexBufferArrayStride.spec.js

@@ -31,7 +31,8 @@ function getPipelineDescriptor(device, testValue) {
 
       }]
 
-    }
+    },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
   };
 }
 

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/capability_checks/limits/maxVertexBuffers.spec.js

@@ -18,7 +18,8 @@
     vertex: {
       module,
       buffers
-    }
+    },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
   };
 }
 

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/createView.spec.js

@@ -6,8 +6,10 @@ import { unreachable } from '../../../common/util/util.js';
 import {
   kTextureAspects,
   kTextureDimensions,
+  kTextureUsages,
   kTextureViewDimensions } from
 '../../capability_info.js';
+import { GPUConst } from '../../constants.js';
 import {
   kTextureFormatInfo,
   kAllTextureFormats,
@@ -338,4 +340,74 @@ fn((t) => {
   t.expectValidationError(() => {
     texture.createView();
   }, state === 'invalid');
+});
+
+g.test('texture_view_usage').
+desc(
+  `Test texture view usage (single, combined, inherited) for every texture format and texture usage`
+).
+params((u) =>
+u //
+.combine('format', kAllTextureFormats).
+combine('textureUsage0', kTextureUsages).
+combine('textureUsage1', kTextureUsages).
+filter(({ format, textureUsage0, textureUsage1 }) => {
+  const info = kTextureFormatInfo[format];
+  const textureUsage = textureUsage0 | textureUsage1;
+
+  if (
+  (textureUsage & GPUConst.TextureUsage.RENDER_ATTACHMENT) !== 0 &&
+  info.color &&
+  !info.colorRender)
+  {
+    return false;
+  }
+
+  return true;
+}).
+beginSubcases().
+combine('textureViewUsage0', [0, ...kTextureUsages]).
+combine('textureViewUsage1', [0, ...kTextureUsages])
+).
+beforeAllSubcases((t) => {
+  const { format, textureUsage0, textureUsage1 } = t.params;
+  const info = kTextureFormatInfo[format];
+  const textureUsage = textureUsage0 | textureUsage1;
+  t.skipIfTextureFormatNotSupported(format);
+  t.selectDeviceOrSkipTestCase(info.feature);
+  if (textureUsage & GPUTextureUsage.STORAGE_BINDING) {
+    t.skipIfTextureFormatNotUsableAsStorageTexture(format);
+  }
+}).
+fn((t) => {
+  const { format, textureUsage0, textureUsage1, textureViewUsage0, textureViewUsage1 } = t.params;
+  const info = kTextureFormatInfo[format];
+
+  const size = [info.blockWidth, info.blockHeight, 1];
+  const dimension = '2d';
+  const mipLevelCount = 1;
+  const usage = textureUsage0 | textureUsage1;
+
+  const textureDescriptor = {
+    size,
+    mipLevelCount,
+    dimension,
+    format,
+    usage
+  };
+
+  const texture = t.createTextureTracked(textureDescriptor);
+
+  let success = true;
+
+  const textureViewUsage = textureViewUsage0 | textureViewUsage1;
+
+  // Texture view usage must be a subset of texture usage
+  if ((~usage & textureViewUsage) !== 0) success = false;
+
+  t.expectValidationError(() => {
+    texture.createView({
+      usage: textureViewUsage
+    });
+  }, !success);
 });

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/gpu_external_texture_expiration.spec.js

@@ -225,7 +225,7 @@ fn(async (t) => {
     } else {
       bindGroup = t.device.createBindGroup({
         layout: t.getDefaultBindGroupLayout(),
-        entries: [{ binding: 0, resource: externalTexture }]
+        entries: [{ binding: 0, resource: mayBeTheSameExternalTexture }]
       });
       t.submitCommandBuffer(bindGroup, true);
     }

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/layout_shader_compat.spec.js

@@ -252,7 +252,8 @@ fn((t) => {
             module: t.device.createShaderModule({
               code: vertexShader
             })
-          }
+          },
+          depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
         });
         break;
       }

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/render_pipeline/float32_blendable.spec.js

@@ -0,0 +1,47 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Tests for capabilities added by float32-blendable flag.
+`;import { makeTestGroup } from '../../../../common/framework/test_group.js';
+
+
+import { CreateRenderPipelineValidationTest } from './common.js';
+
+export const g = makeTestGroup(CreateRenderPipelineValidationTest);
+
+const kFloat32Formats = ['r32float', 'rg32float', 'rgba32float'];
+
+g.test('create_render_pipeline').
+desc(
+  `
+Tests that the float32-blendable feature is required to create a render
+pipeline that uses blending with any float32-format attachment.
+`
+).
+params((u) =>
+u.
+combine('isAsync', [false, true]).
+combine('enabled', [true, false]).
+beginSubcases().
+combine('hasBlend', [true, false]).
+combine('format', kFloat32Formats)
+).
+beforeAllSubcases((t) => {
+  if (t.params.enabled) {
+    t.selectDeviceOrSkipTestCase('float32-blendable');
+  }
+}).
+fn((t) => {
+  const { isAsync, enabled, hasBlend, format } = t.params;
+
+  const descriptor = t.getDescriptor({
+    targets: [
+    {
+      format,
+      blend: hasBlend ? { color: {}, alpha: {} } : undefined
+    }]
+
+  });
+
+  t.doCreateRenderPipelineTest(isAsync, enabled || !hasBlend, descriptor);
+});

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/render_pipeline/fragment_state.spec.js

@@ -5,10 +5,12 @@ This test dedicatedly tests validation of GPUFragmentState of createRenderPipeli
 `;import { makeTestGroup } from '../../../../common/framework/test_group.js';
 import { assert, range } from '../../../../common/util/util.js';
 import {
+  IsDualSourceBlendingFactor,
   kBlendFactors,
   kBlendOperations,
   kMaxColorAttachmentsToTest } from
 '../../../capability_info.js';
+import { GPUConst } from '../../../constants.js';
 import {
   kAllTextureFormats,
   kRenderableColorTextureFormats,
@@ -289,11 +291,17 @@ params((u) =>
 u.
 combine('isAsync', [false, true]).
 combine('component', ['color', 'alpha']).
-beginSubcases().
 combine('srcFactor', kBlendFactors).
 combine('dstFactor', kBlendFactors).
+beginSubcases().
 combine('operation', kBlendOperations)
 ).
+beforeAllSubcases((t) => {
+  const { srcFactor, dstFactor } = t.params;
+  if (IsDualSourceBlendingFactor(srcFactor) || IsDualSourceBlendingFactor(dstFactor)) {
+    t.selectDeviceOrSkipTestCase('dual-source-blending');
+  }
+}).
 fn((t) => {
   const { isAsync, component, srcFactor, dstFactor, operation } = t.params;
 
@@ -308,6 +316,13 @@ fn((t) => {
     operation
   };
   const format = 'rgba8unorm';
+  const useDualSourceBlending =
+  IsDualSourceBlendingFactor(srcFactor) || IsDualSourceBlendingFactor(dstFactor);
+  const fragmentShaderCode = getFragmentShaderCodeWithOutput(
+    [{ values, plainType: 'f32', componentCount: 4 }],
+    null,
+    useDualSourceBlending
+  );
 
   const descriptor = t.getDescriptor({
     targets: [
@@ -317,8 +332,9 @@ fn((t) => {
         color: component === 'color' ? blendComponentToTest : defaultBlendComponent,
         alpha: component === 'alpha' ? blendComponentToTest : defaultBlendComponent
       }
-    }]
+    }],
 
+    fragmentShaderCode
   });
 
   if (operation === 'min' || operation === 'max') {
@@ -407,16 +423,16 @@ fn((t) => {
 
 g.test('pipeline_output_targets,blend').
 desc(
-  `On top of requirements from pipeline_output_targets, when blending is enabled and alpha channel is read indicated by any blend factor, an extra requirement is added:
-  - fragment output must be vec4.
+  `On top of requirements from pipeline_output_targets, when blending is enabled and alpha channel
+    is read indicated by any color blend factor, an extra requirement is added:
+      - fragment output must be vec4.
   `
 ).
 params((u) =>
 u.
 combine('isAsync', [false, true]).
 combine('format', ['r8unorm', 'rg8unorm', 'rgba8unorm', 'bgra8unorm']).
-combine('componentCount', [1, 2, 3, 4]).
-beginSubcases()
+combine('componentCount', [1, 2, 3, 4])
 // The default srcFactor and dstFactor are 'one' and 'zero'. Override just one at a time.
 .combineWithParams([
 ...u.combine('colorSrcFactor', kBlendFactors),
@@ -426,9 +442,21 @@ beginSubcases()
 )
 ).
 beforeAllSubcases((t) => {
-  const { format } = t.params;
+  const { format, colorSrcFactor, colorDstFactor, alphaSrcFactor, alphaDstFactor } = t.params;
+
   const info = kTextureFormatInfo[format];
-  t.selectDeviceOrSkipTestCase(info.feature);
+  const requiredFeatures = [info.feature];
+
+  if (
+  IsDualSourceBlendingFactor(colorSrcFactor) ||
+  IsDualSourceBlendingFactor(colorDstFactor) ||
+  IsDualSourceBlendingFactor(alphaSrcFactor) ||
+  IsDualSourceBlendingFactor(alphaDstFactor))
+  {
+    requiredFeatures.push('dual-source-blending');
+  }
+
+  t.selectDeviceOrSkipTestCase(requiredFeatures);
 }).
 fn((t) => {
   const sampleType = 'float';
@@ -443,6 +471,12 @@ fn((t) => {
   } = t.params;
   const info = kTextureFormatInfo[format];
 
+  const useDualSourceBlending =
+  IsDualSourceBlendingFactor(colorSrcFactor) ||
+  IsDualSourceBlendingFactor(colorDstFactor) ||
+  IsDualSourceBlendingFactor(alphaSrcFactor) ||
+  IsDualSourceBlendingFactor(alphaDstFactor);
+
   const descriptor = t.getDescriptor({
     targets: [
     {
@@ -453,17 +487,147 @@ fn((t) => {
       }
     }],
 
-    fragmentShaderCode: getFragmentShaderCodeWithOutput([
-    { values, plainType: getPlainTypeInfo(sampleType), componentCount }]
+    fragmentShaderCode: getFragmentShaderCodeWithOutput(
+      [{ values, plainType: getPlainTypeInfo(sampleType), componentCount }],
+      null,
+      useDualSourceBlending
     )
   });
 
   const colorBlendReadsSrcAlpha =
-  colorSrcFactor?.includes('src-alpha') || colorDstFactor?.includes('src-alpha');
+  colorSrcFactor?.includes('src-alpha') ||
+  colorDstFactor?.includes('src-alpha') ||
+  colorSrcFactor?.includes('src1-alpha') ||
+  colorDstFactor?.includes('src1-alpha');
   const meetsExtraBlendingRequirement = !colorBlendReadsSrcAlpha || componentCount === 4;
   const _success =
   info.color.type === sampleType &&
   componentCount >= kTexelRepresentationInfo[format].componentOrder.length &&
   meetsExtraBlendingRequirement;
   t.doCreateRenderPipelineTest(isAsync, _success, descriptor);
+});
+
+const kDualSourceBlendingFactors = [
+'src1',
+'one-minus-src1',
+'src1-alpha',
+'one-minus-src1-alpha'];
+
+
+g.test('dual_source_blending,color_target_count').
+desc(
+  `Test that when the blend factor of color attachment 0 uses src1 (the second input of the
+   corresponding blending unit), there must be exactly one color target.
+`
+).
+beforeAllSubcases((t) => t.selectDeviceOrSkipTestCase('dual-source-blending')).
+params((u) =>
+u.
+combine('blendFactor', kDualSourceBlendingFactors).
+combine('colorTargetsCount', [1, 2]).
+combine('maskOutNonZeroIndexColorTargets', [true, false]).
+beginSubcases().
+combine('component', ['color', 'alpha'])
+).
+fn((t) => {
+  const { blendFactor, colorTargetsCount, maskOutNonZeroIndexColorTargets, component } = t.params;
+
+  const defaultBlendComponent = {
+    srcFactor: 'src-alpha',
+    dstFactor: 'dst-alpha',
+    operation: 'add'
+  };
+  const testBlendComponent = {
+    srcFactor: blendFactor,
+    dstFactor: blendFactor,
+    operation: 'add'
+  };
+
+  assert(colorTargetsCount >= 1);
+  const colorTargetStates = new Array(colorTargetsCount);
+  colorTargetStates[0] = {
+    format: 'rgba8unorm',
+    blend: {
+      color: component === 'color' ? testBlendComponent : defaultBlendComponent,
+      alpha: component === 'alpha' ? testBlendComponent : defaultBlendComponent
+    }
+  };
+
+  for (let i = 1; i < colorTargetsCount; ++i) {
+    colorTargetStates[i] = {
+      format: 'rgba8unorm',
+      blend: {
+        color: defaultBlendComponent,
+        alpha: defaultBlendComponent
+      },
+      writeMask: maskOutNonZeroIndexColorTargets ? 0 : GPUConst.ColorWrite.ALL
+    };
+  }
+
+  const descriptor = t.getDescriptor({
+    targets: colorTargetStates,
+    fragmentShaderCode: getFragmentShaderCodeWithOutput(
+      [{ values, plainType: 'f32', componentCount: 4 }],
+      null,
+      true
+    )
+  });
+
+  const isAsync = false;
+  const _success = colorTargetsCount === 1;
+  t.doCreateRenderPipelineTest(isAsync, _success, descriptor);
+});
+
+g.test('dual_source_blending,use_blend_src').
+desc(
+  `Test that when the blend factor of color attachment 0 uses src1, dual source blending must be
+    used in the fragment shader, whether the corresponding color write mask is 0 or not. In
+    contrast, when dual source blending is used in the fragment shader, we don't require blend
+    factor must use src1 (the second input of the corresponding blending unit).
+`
+).
+beforeAllSubcases((t) => t.selectDeviceOrSkipTestCase('dual-source-blending')).
+params((u) =>
+u.
+combine('blendFactor', kBlendFactors).
+combine('useBlendSrc1', [true, false]).
+combine('writeMask', [0, GPUConst.ColorWrite.ALL]).
+beginSubcases().
+combine('component', ['color', 'alpha'])
+).
+fn((t) => {
+  const { blendFactor, useBlendSrc1, writeMask, component } = t.params;
+
+  const defaultBlendComponent = {
+    srcFactor: 'src-alpha',
+    dstFactor: 'dst-alpha',
+    operation: 'add'
+  };
+  const testBlendComponent = {
+    srcFactor: blendFactor,
+    dstFactor: blendFactor,
+    operation: 'add'
+  };
+
+  const descriptor = t.getDescriptor({
+    targets: [
+    {
+      format: 'rgba8unorm',
+      blend: {
+        color: component === 'color' ? testBlendComponent : defaultBlendComponent,
+        alpha: component === 'alpha' ? testBlendComponent : defaultBlendComponent
+      },
+      writeMask
+    }],
+
+    fragmentShaderCode: getFragmentShaderCodeWithOutput(
+      [{ values, plainType: 'f32', componentCount: 4 }],
+      null,
+      useBlendSrc1
+    )
+  });
+
+  const _success = !IsDualSourceBlendingFactor(blendFactor) || useBlendSrc1;
+  const isAsync = false;
+  t.doCreateRenderPipelineTest(isAsync, _success, descriptor);
 });

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/render_pipeline/inter_stage.spec.js

@@ -273,39 +273,29 @@ fn((t) => {
   t.doCreateRenderPipelineTest(isAsync, location < maxInterStageShaderVariables, descriptor);
 });
 
-g.test('max_components_count,output').
+g.test('max_variables_count,output').
 desc(
-  `Tests that validation should fail when scalar components of all user-defined outputs > max vertex shader output components.`
+  `Tests that validation should fail when all user-defined outputs > max vertex shader output
+    variables.`
 ).
 params((u) =>
 u.combine('isAsync', [false, true]).combineWithParams([
-// Number of user-defined output scalar components in test shader =
-//     Math.floor((device.limits.maxInterStageShaderComponents + numScalarDelta) / 4) * 4.
-{ numScalarDelta: 0, topology: 'triangle-list', _success: true },
-{ numScalarDelta: 1, topology: 'triangle-list', _success: false },
-{ numScalarDelta: 0, topology: 'point-list', _success: false },
-{ numScalarDelta: -1, topology: 'point-list', _success: false },
-{ numScalarDelta: -3, topology: 'point-list', _success: false },
-{ numScalarDelta: -4, topology: 'point-list', _success: true }]
+// Number of user-defined output variables in test shader =
+//     device.limits.maxInterStageShaderVariables + numVariablesDelta
+{ numVariablesDelta: 0, topology: 'triangle-list', _success: true },
+{ numVariablesDelta: 1, topology: 'triangle-list', _success: false },
+{ numVariablesDelta: 0, topology: 'point-list', _success: false },
+{ numVariablesDelta: -1, topology: 'point-list', _success: true }]
 )
 ).
 fn((t) => {
-  const { isAsync, numScalarDelta, topology, _success } = t.params;
+  const { isAsync, numVariablesDelta, topology, _success } = t.params;
 
-  const numScalarComponents = t.device.limits.maxInterStageShaderComponents + numScalarDelta;
-
-  const numVec4 = Math.floor(numScalarComponents / 4);
-  const numTrailingScalars = numScalarComponents % 4;
+  const numVec4 = t.device.limits.maxInterStageShaderVariables + numVariablesDelta;
 
   const outputs = range(numVec4, (i) => `@location(${i}) vout${i}: vec4<f32>`);
   const inputs = range(numVec4, (i) => `@location(${i}) fin${i}: vec4<f32>`);
 
-  if (numTrailingScalars > 0) {
-    const typeString = numTrailingScalars === 1 ? 'f32' : `vec${numTrailingScalars}<f32>`;
-    outputs.push(`@location(${numVec4}) vout${numVec4}: ${typeString}`);
-    inputs.push(`@location(${numVec4}) fin${numVec4}: ${typeString}`);
-  }
-
   const descriptor = t.getDescriptorWithStates(
     t.getVertexStateWithOutputs(outputs),
     t.getFragmentStateWithInputs(inputs)
@@ -315,42 +305,32 @@ fn((t) => {
   t.doCreateRenderPipelineTest(isAsync, _success, descriptor);
 });
 
-g.test('max_components_count,input').
+g.test('max_variables_count,input').
 desc(
-  `Tests that validation should fail when scalar components of all user-defined inputs > max vertex shader output components.`
+  `Tests that validation should fail when all user-defined inputs > max vertex shader output
+    variables.`
 ).
 params((u) =>
 u.combine('isAsync', [false, true]).combineWithParams([
-// Number of user-defined input scalar components in test shader =
-//     Math.floor((device.limits.maxInterStageShaderComponents + numScalarDelta) / 4) * 4.
-{ numScalarDelta: 0, useExtraBuiltinInputs: false },
-{ numScalarDelta: 1, useExtraBuiltinInputs: false },
-{ numScalarDelta: 0, useExtraBuiltinInputs: true },
-{ numScalarDelta: -3, useExtraBuiltinInputs: true },
-{ numScalarDelta: -4, useExtraBuiltinInputs: true }]
+// Number of user-defined output variables in test shader =
+//     device.limits.maxInterStageShaderVariables + numVariablesDelta
+{ numVariablesDelta: 0, useExtraBuiltinInputs: false },
+{ numVariablesDelta: 1, useExtraBuiltinInputs: false },
+{ numVariablesDelta: 0, useExtraBuiltinInputs: true },
+{ numVariablesDelta: -1, useExtraBuiltinInputs: true }]
 )
 ).
 fn((t) => {
-  const { isAsync, numScalarDelta, useExtraBuiltinInputs } = t.params;
+  const { isAsync, numVariablesDelta, useExtraBuiltinInputs } = t.params;
 
-  const numScalarComponents =
-  Math.floor((t.device.limits.maxInterStageShaderComponents + numScalarDelta) / 4) * 4;
-  const numExtraComponents = useExtraBuiltinInputs ? t.isCompatibility ? 2 : 3 : 0;
-  const numUsedComponents = numScalarComponents + numExtraComponents;
-  const success = numUsedComponents <= t.device.limits.maxInterStageShaderComponents;
-
-  const numVec4 = Math.floor(numScalarComponents / 4);
-  const numTrailingScalars = numScalarComponents % 4;
+  const numVec4 = t.device.limits.maxInterStageShaderVariables + numVariablesDelta;
+  const numExtraVariables = useExtraBuiltinInputs ? 1 : 0;
+  const numUsedVariables = numVec4 + numExtraVariables;
+  const success = numUsedVariables <= t.device.limits.maxInterStageShaderVariables;
 
   const outputs = range(numVec4, (i) => `@location(${i}) vout${i}: vec4<f32>`);
   const inputs = range(numVec4, (i) => `@location(${i}) fin${i}: vec4<f32>`);
 
-  if (numTrailingScalars > 0) {
-    const typeString = numTrailingScalars === 1 ? 'f32' : `vec${numTrailingScalars}<f32>`;
-    outputs.push(`@location(${numVec4}) vout${numVec4}: ${typeString}`);
-    inputs.push(`@location(${numVec4}) fin${numVec4}: ${typeString}`);
-  }
-
   if (useExtraBuiltinInputs) {
     inputs.push('@builtin(front_facing) front_facing_in: bool');
     if (!t.isCompatibility) {

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/render_pipeline/misc.spec.js

@@ -36,7 +36,7 @@ fn((t) => {
 g.test('vertex_state_only').
 desc(
   `Tests creating vertex-state-only render pipeline. A vertex-only render pipeline has no fragment
-state (and thus has no color state), and can be created with or without depth stencil state.`
+state (and thus has no color state), and must have a depth-stencil state as an attachment is required.`
 ).
 params((u) =>
 u.
@@ -76,7 +76,7 @@ fn((t) => {
     targets: hasColor ? [{ format: 'rgba8unorm' }] : []
   });
 
-  t.doCreateRenderPipelineTest(isAsync, true, descriptor);
+  t.doCreateRenderPipelineTest(isAsync, depthStencilState !== undefined, descriptor);
 });
 
 g.test('pipeline_layout,device_mismatch').

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/resource_usages/texture/in_render_common.spec.js

@@ -179,12 +179,22 @@ combineWithParams([
 combineWithParams([
 { bgLayer: 0, bgLayerCount: 1 },
 { bgLayer: 1, bgLayerCount: 1 },
-{ bgLayer: 1, bgLayerCount: 2 }]
+{ bgLayer: 1, bgLayerCount: 2 },
+{ bgLayer: 0, bgLayerCount: kTextureLayers }]
 ).
 combine('bgUsage', kTextureBindingTypes).
 unless((t) => t.bgUsage !== 'sampled-texture' && t.bgLevelCount > 1).
 combine('inSamePass', [true, false])
 ).
+beforeAllSubcases((t) => {
+  if (t.isCompatibility) {
+    t.skipIf(t.params.bgLayer !== 0, 'view base array layer must equal 0 in compatibility mode');
+    t.skipIf(
+      t.params.bgLayerCount !== kTextureLayers,
+      'view array layers must equal texture array layers in compatibility mode'
+    );
+  }
+}).
 fn((t) => {
   const {
     colorAttachmentLevel,
@@ -287,7 +297,8 @@ combineWithParams([
 combineWithParams([
 { bgLayer: 0, bgLayerCount: 1 },
 { bgLayer: 1, bgLayerCount: 1 },
-{ bgLayer: 1, bgLayerCount: 2 }]
+{ bgLayer: 1, bgLayerCount: 2 },
+{ bgLayer: 0, bgLayerCount: kTextureLayers }]
 ).
 beginSubcases().
 combine('depthReadOnly', [true, false]).
@@ -295,6 +306,15 @@ combine('stencilReadOnly', [true, false]).
 combine('bgAspect', ['depth-only', 'stencil-only']).
 combine('inSamePass', [true, false])
 ).
+beforeAllSubcases((t) => {
+  if (t.isCompatibility) {
+    t.skipIf(t.params.bgLayer !== 0, 'view base array layer must equal 0 in compatibility mode');
+    t.skipIf(
+      t.params.bgLayerCount !== kTextureLayers,
+      'view array layers must equal texture array layers in compatibility mode'
+    );
+  }
+}).
 fn((t) => {
   const {
     dsLevel,
@@ -410,7 +430,8 @@ combine('bg0Levels', [
 combine('bg0Layers', [
 { base: 0, count: 1 },
 { base: 1, count: 1 },
-{ base: 1, count: 2 }]
+{ base: 1, count: 2 },
+{ base: 0, count: kTextureLayers }]
 ).
 combine('bg1Levels', [
 { base: 0, count: 1 },
@@ -420,7 +441,8 @@ combine('bg1Levels', [
 combine('bg1Layers', [
 { base: 0, count: 1 },
 { base: 1, count: 1 },
-{ base: 1, count: 2 }]
+{ base: 1, count: 2 },
+{ base: 0, count: kTextureLayers }]
 ).
 combine('bgUsage0', kTextureBindingTypes).
 combine('bgUsage1', kTextureBindingTypes).
@@ -432,6 +454,18 @@ unless(
 beginSubcases().
 combine('inSamePass', [true, false])
 ).
+beforeAllSubcases((t) => {
+  if (t.isCompatibility) {
+    t.skipIf(
+      t.params.bg0Layers.base !== 0 || t.params.bg1Layers.base !== 0,
+      'view base array layer must equal 0 in compatibility mode'
+    );
+    t.skipIf(
+      t.params.bg0Layers.count !== kTextureLayers || t.params.bg1Layers.count !== kTextureLayers,
+      'view array layers must equal texture array layers in compatibility mode'
+    );
+  }
+}).
 fn((t) => {
   const { bg0Levels, bg0Layers, bg1Levels, bg1Layers, bgUsage0, bgUsage1, inSamePass } = t.params;
 
@@ -523,7 +557,8 @@ combine('view0Levels', [
 combine('view0Layers', [
 { base: 0, count: 1 },
 { base: 1, count: 1 },
-{ base: 1, count: 2 }]
+{ base: 1, count: 2 },
+{ base: 0, count: kTextureLayers }]
 ).
 combine('view1Levels', [
 { base: 0, count: 1 },
@@ -533,12 +568,26 @@ combine('view1Levels', [
 combine('view1Layers', [
 { base: 0, count: 1 },
 { base: 1, count: 1 },
-{ base: 1, count: 2 }]
+{ base: 1, count: 2 },
+{ base: 0, count: kTextureLayers }]
 ).
 combine('aspect0', ['depth-only', 'stencil-only']).
 combine('aspect1', ['depth-only', 'stencil-only']).
 combine('inSamePass', [true, false])
 ).
+beforeAllSubcases((t) => {
+  if (t.isCompatibility) {
+    t.skipIf(
+      t.params.view0Layers.base !== 0 || t.params.view1Layers.base !== 0,
+      'view base array layer must equal 0 in compatibility mode'
+    );
+    t.skipIf(
+      t.params.view0Layers.count !== kTextureLayers ||
+      t.params.view1Layers.count !== kTextureLayers,
+      'view array layers must equal texture array layers in compatibility mode'
+    );
+  }
+}).
 fn((t) => {
   const { view0Levels, view0Layers, view1Levels, view1Layers, aspect0, aspect1, inSamePass } =
   t.params;

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/resource_usages/texture/in_render_misc.spec.js

@@ -4,6 +4,7 @@
 Texture Usages Validation Tests on All Kinds of WebGPU Subresource Usage Scopes.
 `;import { makeTestGroup } from '../../../../../common/framework/test_group.js';
 import { unreachable } from '../../../../../common/util/util.js';
+import { kTextureUsages } from '../../../../capability_info.js';
 import { ValidationTest } from '../../validation_test.js';
 import {
 
@@ -570,4 +571,80 @@ fn((t) => {
   t.expectValidationError(() => {
     encoder.finish();
   }, false);
+});
+
+g.test('subresources,texture_view_usages').
+desc(
+  `
+  Test that the usages of the texture view are used to validate compatibility in command encoding
+  instead of the usages of the base texture.`
+).
+params((u) =>
+u.
+combine('bindingType', ['color-attachment', ...kTextureBindingTypes]).
+combine('viewUsage', [0, ...kTextureUsages])
+).
+fn((t) => {
+  const { bindingType, viewUsage } = t.params;
+
+  const texture = t.createTextureTracked({
+    format: 'r32float',
+    usage:
+    GPUTextureUsage.COPY_SRC |
+    GPUTextureUsage.COPY_DST |
+    GPUTextureUsage.TEXTURE_BINDING |
+    GPUTextureUsage.STORAGE_BINDING |
+    GPUTextureUsage.RENDER_ATTACHMENT,
+    size: [kTextureSize, kTextureSize, 1],
+    ...(t.isCompatibility && {
+      textureBindingViewDimension: '2d-array'
+    })
+  });
+
+  switch (bindingType) {
+    case 'color-attachment':{
+        const encoder = t.device.createCommandEncoder();
+        const renderPassEncoder = encoder.beginRenderPass({
+          colorAttachments: [
+          { view: texture.createView({ usage: viewUsage }), loadOp: 'load', storeOp: 'store' }]
+
+        });
+        renderPassEncoder.end();
+
+        const success = viewUsage === 0 || (viewUsage & GPUTextureUsage.RENDER_ATTACHMENT) !== 0;
+
+        t.expectValidationError(() => {
+          encoder.finish();
+        }, !success);
+        break;
+      }
+    case 'sampled-texture':
+    case 'readonly-storage-texture':
+    case 'writeonly-storage-texture':
+    case 'readwrite-storage-texture':
+      {
+        let success = true;
+        if (viewUsage !== 0) {
+          if (bindingType === 'sampled-texture') {
+            if ((viewUsage & GPUTextureUsage.TEXTURE_BINDING) === 0) success = false;
+          } else {
+            if ((viewUsage & GPUTextureUsage.STORAGE_BINDING) === 0) success = false;
+          }
+        }
+
+        t.expectValidationError(() => {
+          t.createBindGroupForTest(
+            texture.createView({
+              dimension: '2d-array',
+              usage: viewUsage
+            }),
+            bindingType,
+            'unfilterable-float'
+          );
+        }, !success);
+      }
+      break;
+    default:
+      unreachable();
+  }
 });

tests/wpt/webgpu/tests/webgpu/webgpu/api/validation/shader_module/entry_point.spec.js

@@ -127,7 +127,8 @@ fn((t) => {
     vertex: {
       module: t.device.createShaderModule({ code }),
       entryPoint
-    }
+    },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
   };
   let _success = true;
   if (shaderModuleStage !== 'vertex') {
@@ -257,7 +258,8 @@ fn((t) => {
         code
       }),
       entryPoint: undefined
-    }
+    },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
   };
 
   const success = extraShaderModuleStage !== 'vertex';

tests/wpt/webgpu/tests/webgpu/webgpu/capability_info.js

@@ -658,12 +658,29 @@ export const kBlendFactors = [
 'one-minus-dst-alpha',
 'src-alpha-saturated',
 'constant',
-'one-minus-constant'];
+'one-minus-constant',
+'src1',
+'one-minus-src1',
+'src1-alpha',
+'one-minus-src1-alpha'];
 
 
+/** Check if `blendFactor` belongs to the blend factors in the extension "dual-source-blending". */
+export function IsDualSourceBlendingFactor(blendFactor) {
+  switch (blendFactor) {
+    case 'src1':
+    case 'one-minus-src1':
+    case 'src1-alpha':
+    case 'one-minus-src1-alpha':
+      return true;
+    default:
+      return false;
+  }
+}
+
 /** List of all GPUBlendOperation values. */
 export const kBlendOperations = [
-'add', //
+'add',
 'subtract',
 'reverse-subtract',
 'min',
@@ -713,7 +730,6 @@ const [kLimitInfoKeys, kLimitInfoDefaults, kLimitInfoData] =
   'maxBufferSize': [, 268435456, 268435456, kMaxUnsignedLongLongValue],
   'maxVertexAttributes': [, 16, 16],
   'maxVertexBufferArrayStride': [, 2048, 2048],
-  'maxInterStageShaderComponents': [, 64, 60],
   'maxInterStageShaderVariables': [, 16, 15],
 
   'maxColorAttachments': [, 8, 4],
@@ -805,13 +821,16 @@ export const kFeatureNameInfo =
   'depth-clip-control': {},
   'depth32float-stencil8': {},
   'texture-compression-bc': {},
+  'texture-compression-bc-sliced-3d': {},
   'texture-compression-etc2': {},
   'texture-compression-astc': {},
+  'texture-compression-astc-sliced-3d': {},
   'timestamp-query': {},
   'indirect-first-instance': {},
   'shader-f16': {},
   'rg11b10ufloat-renderable': {},
   'float32-filterable': {},
+  'float32-blendable': {},
   'clip-distances': {},
   'dual-source-blending': {}
 };

tests/wpt/webgpu/tests/webgpu/webgpu/compat/api/validation/render_pipeline/unsupported_wgsl.spec.js

@@ -0,0 +1,311 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Tests limitations of createShaderModule in compat mode.
+`;import { makeTestGroup } from '../../../../../common/framework/test_group.js';
+import { keysOf } from '../../../../../common/util/data_tables.js';
+import { kCompatModeUnsupportedStorageTextureFormats } from '../../../../format_info.js';
+import { CompatibilityTest } from '../../../compatibility_test.js';
+
+export const g = makeTestGroup(CompatibilityTest);
+
+g.test('sample_mask').
+desc(
+  `
+Tests that you can not create a render pipeline that uses sample_mask in compat mode.
+
+- Test that a pipeline with a shader that uses sample_mask fails.
+- Test that a pipeline that references a module that has a shader that uses sample_mask
+  but the pipeline does not reference that shader succeeds.
+    `
+).
+params((u) =>
+u.
+combine('entryPoint', ['fsWithoutSampleMaskUsage', 'fsWithSampleMaskUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { entryPoint, async } = t.params;
+  const module = t.device.createShaderModule({
+    code: `
+        @vertex fn vs() -> @builtin(position) vec4f {
+            return vec4f(1);
+        }
+        struct Output {
+            @builtin(sample_mask) mask_out: u32,
+            @location(0) color : vec4f,
+        }
+        @fragment fn fsWithoutSampleMaskUsage() -> @location(0) vec4f {
+            return vec4f(1.0, 1.0, 1.0, 1.0);
+        }
+        @fragment fn fsWithSampleMaskUsage() -> Output {
+            var o: Output;
+            // We need to make sure this sample_mask isn't optimized out even if its value equals "no op".
+            o.mask_out = 0xFFFFFFFFu;
+            o.color = vec4f(1.0, 1.0, 1.0, 1.0);
+            return o;
+        }
+      `
+  });
+  const isValid = !t.isCompatibility || entryPoint === 'fsWithoutSampleMaskUsage';
+  t.doCreateRenderPipelineTest(async, isValid, {
+    layout: 'auto',
+    vertex: { module },
+    fragment: {
+      module,
+      entryPoint,
+      targets: [{ format: 'rgba8unorm' }]
+    }
+  });
+});
+
+g.test('sample_index').
+desc(
+  `
+Tests that you can not create a render pipeline that uses sample_index in compat mode.
+
+- Test that a pipeline with a shader that uses sample_index fails.
+- Test that a pipeline that references a module that has a shader that uses sample_index
+  but the pipeline does not reference that shader succeeds.
+    `
+).
+params((u) =>
+u.
+combine('entryPoint', ['fsWithoutSampleIndexUsage', 'fsWithSampleIndexUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { entryPoint, async } = t.params;
+
+  const module = t.device.createShaderModule({
+    code: `
+        @vertex fn vs() -> @builtin(position) vec4f {
+            return vec4f(1);
+        }
+        @fragment fn fsWithoutSampleIndexUsage() -> @location(0) vec4f {
+            return vec4f(0);
+        }
+        @fragment fn fsWithSampleIndexUsage(@builtin(sample_index) sampleIndex: u32) -> @location(0) vec4f {
+            _ = sampleIndex;
+            return vec4f(0);
+        }
+          `
+  });
+
+  const isValid = !t.isCompatibility || entryPoint === 'fsWithoutSampleIndexUsage';
+  t.doCreateRenderPipelineTest(async, isValid, {
+    layout: 'auto',
+    vertex: { module },
+    fragment: {
+      module,
+      entryPoint,
+      targets: [{ format: 'rgba8unorm' }]
+    }
+  });
+});
+
+g.test('interpolate').
+desc(
+  `Tests that you can not create a render pipeline that uses interpolate(linear), interpolate(...,sample),
+     interpolate(flat), nor interpolate(flat, first) in compat mode.`
+).
+params((u) =>
+u.
+combineWithParams([
+{ success: false, interpolate: '@interpolate(linear)' },
+{ success: false, interpolate: '@interpolate(linear, sample)' },
+{ success: false, interpolate: '@interpolate(perspective, sample)' },
+{ success: false, interpolate: '@interpolate(flat)' },
+{ success: false, interpolate: '@interpolate(flat, first)' },
+{ success: true, interpolate: '@interpolate(flat, either)' }]
+).
+combine('entryPoint', [
+'fsWithoutInterpolationUsage',
+'fsWithInterpolationUsage1',
+'fsWithInterpolationUsage2',
+'fsWithInterpolationUsage3']
+).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { interpolate, success, entryPoint, async } = t.params;
+
+  const module = t.device.createShaderModule({
+    code: `
+        struct Vertex {
+            @builtin(position) pos: vec4f,
+            @location(0) ${interpolate} color : vec4f,
+        };
+        @vertex fn vs() -> Vertex {
+            var v: Vertex;
+            v.pos = vec4f(1);
+            v.color = vec4f(1);
+            return v;
+        }
+        @fragment fn fsWithoutInterpolationUsage() -> @location(0) vec4f {
+            return vec4f(1);
+        }
+        @fragment fn fsWithInterpolationUsage1(v: Vertex) -> @location(0) vec4f {
+            return vec4f(1);
+        }
+        @fragment fn fsWithInterpolationUsage2(v: Vertex) -> @location(0) vec4f {
+            return v.pos;
+        }
+        @fragment fn fsWithInterpolationUsage3(v: Vertex) -> @location(0) vec4f {
+            return v.color;
+        }
+        `
+  });
+
+  const isValid = success || !t.isCompatibility || entryPoint === 'fsWithoutInterpolationUsage';
+  t.doCreateRenderPipelineTest(async, isValid, {
+    layout: 'auto',
+    vertex: { module },
+    fragment: {
+      entryPoint,
+      module,
+      targets: [{ format: 'rgba8unorm' }]
+    }
+  });
+});
+
+g.test('unsupportedStorageTextureFormats,computePipeline').
+desc(
+  `Tests that you can not create a compute pipeline that uses an
+     unsupported storage texture format in compat mode.`
+).
+params((u) =>
+u.
+combine('format', kCompatModeUnsupportedStorageTextureFormats).
+combine('entryPoint', ['csWithoutStorageUsage', 'csWithStorageUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { format, entryPoint, 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 csWithoutStorageUsage() {
+        }
+        @compute @workgroup_size(1) fn csWithStorageUsage() {
+            _ = textureLoad(s, vec2u(0));
+        }
+      `
+  });
+
+  const isValid = !t.isCompatibility || entryPoint === 'csWithoutStorageUsage';
+  t.doCreateComputePipelineTest(async, isValid, {
+    layout: 'auto',
+    compute: { module, entryPoint }
+  });
+});
+
+g.test('unsupportedStorageTextureFormats,renderPipeline').
+desc(
+  `Tests that you can not create a render pipeline that uses an
+     unsupported storage texture format in compat mode.`
+).
+params((u) =>
+u.
+combine('format', kCompatModeUnsupportedStorageTextureFormats).
+combine('entryPoint', ['vsWithoutStorageUsage', 'vsWithStorageUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { format, entryPoint, async } = t.params;
+
+  const module = t.device.createShaderModule({
+    code: `
+        @group(0) @binding(0) var s: texture_storage_2d<${format}, read>;
+        @vertex fn vsWithoutStorageUsage() -> @builtin(position) vec4f {
+            return vec4f(0);
+        }
+        @vertex fn vsWithStorageUsage() -> @builtin(position) vec4f {
+            _ = textureLoad(s, vec2u(0));
+            return vec4f(0);
+        }
+      `
+  });
+
+  const isValid = !t.isCompatibility || entryPoint === 'vsWithoutStorageUsage';
+  t.doCreateRenderPipelineTest(async, isValid, {
+    layout: 'auto',
+    vertex: { module, entryPoint },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
+  });
+});
+
+const kDepthTextureTypeToParams = {
+  texture_depth_2d: 'vec2u(0), 0',
+  texture_depth_2d_array: 'vec2u(0), 0, 0',
+  texture_depth_multisampled_2d: 'vec2u(0), 0'
+};
+const kDepthTextureTypes = keysOf(kDepthTextureTypeToParams);
+
+g.test('textureLoad_with_depth_textures,computePipeline').
+desc(
+  `Tests that you can not create a compute pipeline that uses textureLoad with a depth texture in compat mode.`
+).
+params((u) =>
+u.
+combine('type', kDepthTextureTypes).
+combine('entryPoint', ['csWithoutDepthUsage', 'csWithDepthUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { type, entryPoint, async } = t.params;
+  const params = kDepthTextureTypeToParams[type];
+
+  const module = t.device.createShaderModule({
+    code: `
+        @group(0) @binding(0) var t: ${type};
+        @compute @workgroup_size(1) fn csWithoutDepthUsage() {
+        }
+        @compute @workgroup_size(1) fn csWithDepthUsage() {
+            _ = textureLoad(t, ${params});
+        }
+      `
+  });
+
+  const isValid = !t.isCompatibility || entryPoint === 'csWithoutDepthUsage';
+  t.doCreateComputePipelineTest(async, isValid, {
+    layout: 'auto',
+    compute: { module, entryPoint }
+  });
+});
+
+g.test('textureLoad_with_depth_textures,renderPipeline').
+desc(
+  `Tests that you can not create a render pipeline that uses textureLoad with a depth texture in compat mode.`
+).
+params((u) =>
+u.
+combine('type', kDepthTextureTypes).
+combine('entryPoint', ['vsWithoutDepthUsage', 'vsWithDepthUsage']).
+combine('async', [false, true])
+).
+fn((t) => {
+  const { type, entryPoint, async } = t.params;
+  const params = kDepthTextureTypeToParams[type];
+
+  const module = t.device.createShaderModule({
+    code: `
+        @group(0) @binding(0) var t: ${type};
+        @vertex fn vsWithoutDepthUsage() -> @builtin(position) vec4f {
+            return vec4f(0);
+        }
+        @vertex fn vsWithDepthUsage() -> @builtin(position) vec4f {
+            _ = textureLoad(t, ${params});
+            return vec4f(0);
+        }
+      `
+  });
+
+  const isValid = !t.isCompatibility || entryPoint === 'vsWithoutDepthUsage';
+  t.doCreateRenderPipelineTest(async, isValid, {
+    layout: 'auto',
+    vertex: { module, entryPoint },
+    depthStencil: { format: 'depth32float', depthWriteEnabled: true, depthCompare: 'always' }
+  });
+});

tests/wpt/webgpu/tests/webgpu/webgpu/compat/api/validation/shader_module/shader_module.spec.js

@@ -1,156 +0,0 @@
-/**
-* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
-**/export const description = `
-Tests limitations of createShaderModule in compat mode.
-`;import { makeTestGroup } from '../../../../../common/framework/test_group.js';
-import { keysOf } from '../../../../../common/util/data_tables.js';
-import { kCompatModeUnsupportedStorageTextureFormats } from '../../../../format_info.js';
-import { CompatibilityTest } from '../../../compatibility_test.js';
-
-export const g = makeTestGroup(CompatibilityTest);
-
-g.test('sample_mask').
-desc(`Tests that you can not create a shader module that uses sample_mask in compat mode.`).
-fn((t) => {
-  t.expectGPUErrorInCompatibilityMode(
-    'validation',
-    () =>
-    t.device.createShaderModule({
-      code: `
-            @vertex fn vs() -> @builtin(position) vec4f {
-                return vec4f(1);
-            }
-            struct Output {
-                @builtin(sample_mask) mask_out: u32,
-                @location(0) color : vec4f,
-            }
-            @fragment fn fsWithSampleMaskUsage() -> Output {
-                var o: Output;
-                // We need to make sure this sample_mask isn't optimized out even if its value equals "no op".
-                o.mask_out = 0xFFFFFFFFu;
-                o.color = vec4f(1.0, 1.0, 1.0, 1.0);
-                return o;
-            }
-          `
-    }),
-    true
-  );
-});
-
-g.test('sample_index').
-desc(`Tests that you can not create a shader module that uses sample_index in compat mode.`).
-fn((t) => {
-  t.expectGPUErrorInCompatibilityMode(
-    'validation',
-    () =>
-    t.device.createShaderModule({
-      code: `
-            @vertex fn vs() -> @builtin(position) vec4f {
-                return vec4f(1);
-            }
-            @fragment fn fsWithSampleIndexUsage(@builtin(sample_index) sampleIndex: u32) -> @location(0) vec4f {
-                _ = sampleIndex;
-                return vec4f(0);
-            }
-              `
-    }),
-    true
-  );
-});
-
-g.test('interpolate').
-desc(
-  `Tests that you can not create a shader module that uses interpolate(linear), interpolate(...,sample),
-     interpolate(flat), nor interpolate(flat, first) in compat mode.`
-).
-params((u) =>
-u.combineWithParams([
-{ success: true, interpolate: '' },
-{ success: false, interpolate: '@interpolate(linear)' },
-{ success: false, interpolate: '@interpolate(linear, sample)' },
-{ success: false, interpolate: '@interpolate(perspective, sample)' },
-{ success: false, interpolate: '@interpolate(flat)' },
-{ success: false, interpolate: '@interpolate(flat, first)' },
-{ success: true, interpolate: '@interpolate(flat, either)' }]
-)
-).
-fn((t) => {
-  const { interpolate, success } = t.params;
-
-  t.expectGPUErrorInCompatibilityMode(
-    'validation',
-    () =>
-    t.device.createShaderModule({
-      code: `
-            struct Vertex {
-                @builtin(position) pos: vec4f,
-                @location(0) ${interpolate} color : vec4f,
-            };
-            @vertex fn vs() -> Vertex {
-                var v: Vertex;
-                v.pos = vec4f(1);
-                v.color = vec4f(1);
-                return v;
-            }
-            @fragment fn fsWithInterpolationUsage(v: Vertex) -> @location(0) vec4f {
-                return v.color;
-            }
-            `
-    }),
-    !success
-  );
-});
-
-g.test('unsupportedStorageTextureFormats').
-desc(
-  `Tests that you can not create a shader module with unsupported storage texture formats in compat mode.`
-).
-params((u) => u.combine('format', kCompatModeUnsupportedStorageTextureFormats)).
-fn((t) => {
-  const { format } = t.params;
-
-  t.expectGPUErrorInCompatibilityMode(
-    'validation',
-    () =>
-    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));
-            }
-          `
-    }),
-    true
-  );
-});
-
-const kDepthTextureTypeToParams = {
-  texture_depth_2d: 'vec2u(0), 0',
-  texture_depth_2d_array: 'vec2u(0), 0, 0',
-  texture_depth_multisampled_2d: 'vec2u(0), 0'
-};
-const kDepthTextureTypes = keysOf(kDepthTextureTypeToParams);
-
-g.test('textureLoad_with_depth_textures').
-desc(
-  `Tests that you can not create a shader module that uses textureLoad with a depth texture in compat mode.`
-).
-params((u) => u.combine('type', kDepthTextureTypes)).
-fn((t) => {
-  const { type } = t.params;
-  const params = kDepthTextureTypeToParams[type];
-
-  t.expectGPUErrorInCompatibilityMode(
-    'validation',
-    () =>
-    t.device.createShaderModule({
-      code: `
-            @group(0) @binding(0) var t: ${type};
-            @compute @workgroup_size(1) fn cs() {
-                _ = textureLoad(t, ${params});
-            }
-          `
-    }),
-    true
-  );
-});

tests/wpt/webgpu/tests/webgpu/webgpu/format_info.js

@@ -1785,6 +1785,14 @@ export function isDepthOrStencilTextureFormat(format) {
   return isDepthTextureFormat(format) || isStencilTextureFormat(format);
 }
 
+export function isEncodableTextureFormat(format) {
+  return kEncodableTextureFormats.includes(format);
+}
+
+export function canUseAsRenderTarget(format) {
+  return kTextureFormatInfo[format].colorRender || isDepthOrStencilTextureFormat(format);
+}
+
 export const kCompatModeUnsupportedStorageTextureFormats = [
 'rg32float',
 'rg32sint',
@@ -1815,6 +1823,13 @@ export function isCompressedFloatTextureFormat(format) {
   return isCompressedTextureFormat(format) && format.includes('float');
 }
 
+/**
+ * Returns true of format can be multisampled.
+ */
+export function isMultisampledTextureFormat(format) {
+  return kAllTextureFormatInfo[format].multisample;
+}
+
 export const kFeaturesForFormats = getFeaturesForFormats(kAllTextureFormats);
 
 /**

tests/wpt/webgpu/tests/webgpu/webgpu/listing.js

@@ -638,6 +638,15 @@ export const listing = [
     ],
     "readme": "Test every method or option that shouldn't be allowed without a feature enabled.\nIf the feature is not enabled, any use of an enum value added by a feature must be an\n*exception*, per <https://github.com/gpuweb/gpuweb/blob/main/design/ErrorConventions.md>.\n\n- x= that feature {enabled, disabled}\n\nGenerally one file for each feature name, but some may be grouped (e.g. one file for all optional\nquery types, one file for all optional texture formats).\n\nTODO: implement"
   },
+  {
+    "file": [
+      "api",
+      "validation",
+      "capability_checks",
+      "features",
+      "clip_distances"
+    ]
+  },
   {
     "file": [
       "api",
@@ -791,15 +800,6 @@ export const listing = [
       "maxDynamicUniformBuffersPerPipelineLayout"
     ]
   },
-  {
-    "file": [
-      "api",
-      "validation",
-      "capability_checks",
-      "limits",
-      "maxInterStageShaderComponents"
-    ]
-  },
   {
     "file": [
       "api",
@@ -1435,6 +1435,14 @@ export const listing = [
       "depth_stencil_state"
     ]
   },
+  {
+    "file": [
+      "api",
+      "validation",
+      "render_pipeline",
+      "float32_blendable"
+    ]
+  },
   {
     "file": [
       "api",
@@ -1697,7 +1705,7 @@ export const listing = [
       "api",
       "validation",
       "render_pipeline",
-      "vertex_state"
+      "unsupported_wgsl"
     ]
   },
   {
@@ -1705,8 +1713,8 @@ export const listing = [
       "compat",
       "api",
       "validation",
-      "shader_module",
-      "shader_module"
+      "render_pipeline",
+      "vertex_state"
     ]
   },
   {
@@ -2978,6 +2986,26 @@ export const listing = [
       "pow"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "quadBroadcast"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "quadSwap"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -3128,6 +3156,76 @@ export const listing = [
       "storageBarrier"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupAdd"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupAll"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupAny"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBallot"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBitwise"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBroadcast"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupMul"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -3228,6 +3326,16 @@ export const listing = [
       "textureSample"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "expression",
+      "call",
+      "builtin",
+      "textureSampleBaseClampToEdge"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -3789,6 +3897,14 @@ export const listing = [
       "user_io"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "shader_io",
+      "vertex_builtins"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -3835,6 +3951,14 @@ export const listing = [
       "increment_decrement"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "execution",
+      "statement",
+      "phony"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -4009,6 +4133,15 @@ export const listing = [
       "parse"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "binary",
+      "short_circuiting_and_or"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -4609,6 +4742,26 @@ export const listing = [
       "pow"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "quadBroadcast"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "quadSwap"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -4749,6 +4902,106 @@ export const listing = [
       "step"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupAdd"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupAnyAll"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBallot"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBitwise"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBroadcast"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupBroadcastFirst"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupElect"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupMinMax"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupMul"
+    ]
+  },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "expression",
+      "call",
+      "builtin",
+      "subgroupShuffle"
+    ]
+  },
   {
     "file": [
       "shader",
@@ -5143,6 +5396,14 @@ export const listing = [
       "logical_negation"
     ]
   },
+  {
+    "file": [
+      "shader",
+      "validation",
+      "extension",
+      "clip_distances"
+    ]
+  },
   {
     "file": [
       "shader",

tests/wpt/webgpu/tests/webgpu/webgpu/print_environment.spec.js

@@ -35,9 +35,7 @@ NOTE: If your test runtime elides logs when tests pass, you won't see the prints
 in the logs. On non-WPT runtimes, it will also print to the console with console.log.
 WPT disallows console.log and doesn't support logs on passing tests, so this does nothing on WPT.`
 ).
-fn(async (t) => {
-  // MAINTENANCE_TODO: Remove requestAdapterInfo when info is implemented.
-  const adapterInfo = t.adapter.info || (await t.adapter.requestAdapterInfo());
+fn((t) => {
   const isCompatibilityMode = t.adapter.
   isCompatibilityMode;
 
@@ -51,7 +49,7 @@ fn(async (t) => {
       adapter: {
         isFallbackAdapter: t.adapter.isFallbackAdapter,
         isCompatibilityMode,
-        info: adapterInfo,
+        info: t.adapter.info,
         features: Array.from(t.adapter.features),
         limits: t.adapter.limits
       }

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/access/matrix/index.spec.js

@@ -11,7 +11,9 @@ import {
   abstractFloat,
   f32,
   vec } from
+
 '../../../../../util/conversion.js';
+import { align } from '../../../../../util/math.js';
 
 import { allInputSources, basicExpressionBuilder, run } from '../../expression.js';
 
@@ -197,4 +199,74 @@ fn(async (t) => {
     { inputSource: 'const' },
     cases
   );
+});
+
+g.test('non_const_index').
+specURL('https://www.w3.org/TR/WGSL/#matrix-access-expr').
+desc(`Test indexing of a matrix using non-const index`).
+params((u) => u.combine('columns', [2, 3, 4]).combine('rows', [2, 3, 4])).
+fn((t) => {
+  const cols = t.params.columns;
+  const rows = t.params.rows;
+  const values = Array.from(Array(cols * rows).keys());
+  const wgsl = `
+@group(0) @binding(0) var<storage, read_write> output : array<f32, ${cols * rows}>;
+
+@compute @workgroup_size(${cols}, ${rows})
+fn main(@builtin(local_invocation_id) invocation_id : vec3<u32>) {
+  let m = mat${cols}x${rows}f(${values.join(', ')});
+  output[invocation_id.x*${rows} + invocation_id.y] = m[invocation_id.x][invocation_id.y];
+}
+`;
+
+  const pipeline = t.device.createComputePipeline({
+    layout: 'auto',
+    compute: {
+      module: t.device.createShaderModule({ code: wgsl }),
+      entryPoint: 'main'
+    }
+  });
+
+  const bufferSize = (arr) => {
+    let offset = 0;
+    let alignment = 0;
+    for (const value of arr) {
+      alignment = Math.max(alignment, value.type.alignment);
+      offset = align(offset, value.type.alignment) + value.type.size;
+    }
+    return align(offset, alignment);
+  };
+
+  const toArray = (arr) => {
+    const array = new Uint8Array(bufferSize(arr));
+    let offset = 0;
+    for (const value of arr) {
+      offset = align(offset, value.type.alignment);
+      value.copyTo(array, offset);
+      offset += value.type.size;
+    }
+    return array;
+  };
+
+  const expected = values.map((i) => Type['f32'].create(i));
+
+  const outputBuffer = t.createBufferTracked({
+    size: bufferSize(expected),
+    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
+  });
+
+  const bindGroup = t.device.createBindGroup({
+    layout: pipeline.getBindGroupLayout(0),
+    entries: [{ binding: 0, resource: { buffer: outputBuffer } }]
+  });
+
+  const encoder = t.device.createCommandEncoder();
+  const pass = encoder.beginComputePass();
+  pass.setPipeline(pipeline);
+  pass.setBindGroup(0, bindGroup);
+  pass.dispatchWorkgroups(1);
+  pass.end();
+  t.queue.submit([encoder.finish()]);
+
+  t.expectGPUBufferValuesEqual(outputBuffer, toArray(expected));
 });

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/quadBroadcast.spec.js

@@ -0,0 +1,656 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for quadBroadcast.
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
+import { assert, unreachable } from '../../../../../../common/util/util.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import { kBit } from '../../../../../util/constants.js';
+import {
+  kConcreteNumericScalarsAndVectors,
+  Type,
+  VectorType,
+  scalarTypeOf } from
+'../../../../../util/conversion.js';
+import { align } from '../../../../../util/math.js';
+
+import {
+  kWGSizes,
+  kDataSentinel,
+  kPredicateCases,
+  runComputeTest,
+  SubgroupTest,
+  kFramebufferSizes,
+  runFragmentTest } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(SubgroupTest);
+
+const kTypes = objectsToRecord(kConcreteNumericScalarsAndVectors);
+
+/**
+ * Generates scalar values for type
+ *
+ * Generates 4 32-bit values whose bit patterns represent
+ * interesting values of the data type.
+ * @param type The data type
+ */
+function generateScalarValues(type) {
+  const scalarTy = scalarTypeOf(type);
+  switch (scalarTy) {
+    case Type.u32:
+      return [kBit.u32.min, kBit.u32.max, 1111, 2222];
+    case Type.i32:
+      return [
+      kBit.i32.positive.min,
+      kBit.i32.positive.max,
+      kBit.i32.negative.min,
+      0xffffffff // -1
+      ];
+    case Type.f32:
+      return [
+      kBit.f32.positive.zero,
+      kBit.f32.positive.nearest_max,
+      kBit.f32.negative.nearest_min,
+      0xbf800000 // -1
+      ];
+    case Type.f16:
+      return [
+      kBit.f16.positive.zero,
+      kBit.f16.positive.nearest_max,
+      kBit.f16.negative.nearest_min,
+      0xbc00 // -1
+      ];
+    default:
+      unreachable(`Unsupported type: ${type.toString()}`);
+  }
+  return [0, 0, 0, 0];
+}
+
+/**
+ * Generates input bit patterns for the input type
+ *
+ * Generates 4 values of type in a Uint32Array.
+ * 16-bit types are appropriately packed.
+ * @param type The data type
+ */
+function generateTypedInputs(type) {
+  const scalarValues = generateScalarValues(type);
+  let elements = 1;
+  if (type instanceof VectorType) {
+    elements = type.width;
+  }
+  if (type.requiresF16()) {
+    switch (elements) {
+      case 1:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[3] << 16]
+        );
+      case 2:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[3] | scalarValues[3] << 16]
+        );
+      case 3:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[0] | kDataSentinel << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[1] | kDataSentinel << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[2] | kDataSentinel << 16,
+        scalarValues[3] | scalarValues[3] << 16,
+        scalarValues[3] | kDataSentinel << 16]
+        );
+      case 4:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[3] | scalarValues[3] << 16,
+        scalarValues[3] | scalarValues[3] << 16]
+        );
+      default:
+        unreachable(`Unsupported type: ${type.toString()}`);
+    }
+    return new Uint32Array([0]);
+  } else {
+    const bound = elements === 3 ? 4 : elements;
+    const values = [];
+    for (let i = 0; i < 4; i++) {
+      for (let j = 0; j < bound; j++) {
+        if (j < elements) {
+          values.push(scalarValues[i]);
+        } else {
+          values.push(kDataSentinel);
+        }
+      }
+    }
+    return new Uint32Array(values);
+  }
+}
+
+/**
+ * Checks results from data types test
+ *
+ * The output is expected to match the input values corresponding to the
+ * id being broadcast (assuming a linear mapping).
+ * @param metadata An unused parameter
+ * @param output The output data
+ * @param input The input data
+ * @param broadcast The id being broadcast
+ * @param type The data type being tested
+ */
+function checkDataTypes(
+metadata, // unused
+output,
+input,
+broadcast,
+type)
+{
+  if (type.requiresF16() && !(type instanceof VectorType)) {
+    const expectIdx = Math.floor(broadcast / 2);
+    const expectShift = broadcast % 2 === 1;
+    let expect = input[expectIdx];
+    if (expectShift) {
+      expect >>= 16;
+    }
+    expect &= 0xffff;
+
+    for (let i = 0; i < 4; i++) {
+      const index = Math.floor(i / 2);
+      const shift = i % 2 === 1;
+      let res = output[index];
+      if (shift) {
+        res >>= 16;
+      }
+      res &= 0xffff;
+      if (res !== expect) {
+        return new Error(`${i}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+      }
+    }
+  } else {
+    let uints = 1;
+    if (type instanceof VectorType) {
+      uints = type.width === 3 ? 4 : type.width;
+      if (type.requiresF16()) {
+        uints = Math.floor(uints / 2);
+      }
+    }
+    for (let i = 0; i < 4; i++) {
+      for (let j = 0; j < uints; j++) {
+        const expect = input[broadcast * uints + j];
+        const res = output[i * uints + j];
+        if (res !== expect) {
+          return new Error(`${i * uints + j}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+        }
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('data_types').
+desc('Test allowed data types').
+params((u) =>
+u.
+combine('type', keysOf(kTypes)).
+beginSubcases().
+combine('id', [0, 1, 2, 3])
+).
+beforeAllSubcases((t) => {
+  const features = ['subgroups'];
+  const type = kTypes[t.params.type];
+  if (type.requiresF16()) {
+    features.push('subgroups-f16');
+    features.push('shader-f16');
+  }
+  t.selectDeviceOrSkipTestCase(features);
+}).
+fn(async (t) => {
+  const wgSize = [4, 1, 1];
+  const type = kTypes[t.params.type];
+  let enables = `enable subgroups;\n`;
+  if (type.requiresF16()) {
+    enables += `enable f16;\nenable subgroups_f16;`;
+  }
+  const wgsl = `
+${enables}
+
+@group(0) @binding(0)
+var<storage> input : array<${type.toString()}>;
+
+@group(0) @binding(1)
+var<storage, read_write> output : array<${type.toString()}>;
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : array<u32>; // unused
+
+@compute @workgroup_size(${wgSize[0]}, ${wgSize[1]}, ${wgSize[2]})
+fn main(
+  @builtin(subgroup_invocation_id) id : u32,
+) {
+  // Force usage
+  _ = metadata[0];
+
+  output[id] = quadBroadcast(input[id], ${t.params.id});
+}`;
+
+  const inputData = generateTypedInputs(type);
+  let uintsPerOutput = 1;
+  if (type instanceof VectorType) {
+    uintsPerOutput = type.width === 3 ? 4 : type.width;
+    if (type.requiresF16()) {
+      uintsPerOutput = Math.floor(uintsPerOutput / 2);
+    }
+  }
+  await runComputeTest(
+    t,
+    wgsl,
+    wgSize,
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkDataTypes(metadata, output, inputData, t.params.id, type);
+    }
+  );
+});
+
+/**
+ * Checks quadBroadcast in compute shaders
+ *
+ * Assumes that quads are linear within a subgroup.
+ *
+ * @param metadata An array of integers divided as follows:
+ *                 * first half subgroup invocation ids
+ *                 * second half subgroup sizes
+ * @param output An array of integers divided as follows:
+ *               * first half results of quad broadcast
+ *               * second half generated unique subgroup ids
+ * @param broadcast The id being broadcast in the range [0, 3]
+ * @param filter A functor to filter active invocations
+ */
+function checkBroadcastCompute(
+metadata,
+output,
+broadcast,
+filter)
+{
+  assert(broadcast === Math.trunc(broadcast));
+  assert(broadcast >= 0 && broadcast <= 3);
+
+  const bound = Math.floor(output.length / 2);
+  for (let i = 0; i < bound; i++) {
+    const subgroup_id = output[bound + i];
+    const id = metadata[i];
+    const size = metadata[bound + i];
+    if (!filter(id, size)) {
+      if (output[i] !== kDataSentinel) {
+        return new Error(`Unexpected write for invocation ${i}`);
+      }
+      continue;
+    }
+
+    const quad_id = Math.floor(id / 4);
+    const quad = [-1, -1, -1, -1];
+    for (let j = 0; j < bound; j++) {
+      const other_id = metadata[j];
+      const other_quad_id = Math.floor(other_id / 4);
+      const other_quad_index = other_id % 4;
+      const other_subgroup_id = output[bound + j];
+      if (other_subgroup_id === subgroup_id && quad_id === other_quad_id) {
+        quad[other_quad_index] = j;
+      }
+    }
+    for (let j = 0; j < 4; j++) {
+      if (quad[j] === -1) {
+        return new Error(`Invocation ${i}: missing quad index ${j}`);
+      }
+    }
+    for (let j = 0; j < 4; j++) {
+      if (output[quad[j]] !== output[quad[broadcast]]) {
+        return new Error(`Incorrect result for quad: base invocation = ${
+        quad[broadcast]
+        }, invocation = ${quad[j]}
+- expected: ${output[quad[broadcast]]}
+-      got: ${output[quad[j]]}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('compute,all_active').
+desc(
+  `Tests broadcast with all active invocations
+
+Quad operations require a full quad so workgroup sizes are limited to multiples of 4.
+  `
+).
+params((u) =>
+u.
+combine('wgSize', kWGSizes).
+filter((t) => {
+  const wgThreads = t.wgSize[0] * t.wgSize[1] * t.wgSize[2];
+  return wgThreads % 4 === 0;
+}).
+beginSubcases().
+combine('id', [0, 1, 2, 3])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : u32; // unused
+
+struct Output {
+  results : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(1)
+var<storage, read_write> output : Output;
+
+struct Metadata {
+  id : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  // Force usage
+  _ = inputs;
+
+  let b = quadBroadcast(lid, ${t.params.id});
+  output.results[lid] = b;
+  output.subgroup_size[lid] = subgroupBroadcastFirst(lid + 1);
+  metadata.id[lid] = id;
+  metadata.subgroup_size[lid] = subgroupSize;
+}`;
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    new Uint32Array([0]), // unused
+    (metadata, output) => {
+      return checkBroadcastCompute(metadata, output, t.params.id, (id, size) => {
+        return true;
+      });
+    }
+  );
+});
+
+g.test('compute,split').
+desc(
+  `Tests broadcast with predicated invocations
+
+Quad operations require a full quad so workgroup sizes are limited to multiples of 4.
+Quad operations require a fully active quad to operate correctly so several of the
+predication filters are skipped.
+  `
+).
+params((u) =>
+u.
+combine('predicate', keysOf(kPredicateCases)).
+filter((t) => {
+  return t.predicate === 'lower_half' || t.predicate === 'upper_half';
+}).
+combine('wgSize', kWGSizes).
+filter((t) => {
+  const wgThreads = t.wgSize[0] * t.wgSize[1] * t.wgSize[2];
+  return wgThreads % 4 === 0;
+}).
+beginSubcases().
+combine('id', [0, 1, 2, 3])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+  const testcase = kPredicateCases[t.params.predicate];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : u32; // unused
+
+struct Output {
+  results : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(1)
+var<storage, read_write> output : Output;
+
+struct Metadata {
+  id : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  // Force usage
+  _ = inputs;
+
+  output.subgroup_size[lid] = subgroupBroadcastFirst(lid + 1);
+  metadata.id[lid] = id;
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  if ${testcase.cond} {
+    let b = quadBroadcast(lid, ${t.params.id});
+    output.results[lid] = b;
+  }
+}`;
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    new Uint32Array([0]), // unused
+    (metadata, output) => {
+      return checkBroadcastCompute(metadata, output, t.params.id, testcase.filter);
+    }
+  );
+});
+
+/**
+ * Checks results of quadBroadcast in fragment shaders.
+ *
+ * @param data The framebuffer output
+ *             * component 0 is the broadcast of the integer x position
+ *             * component 1 is the broadcast of the integer y position
+ * @param format The framebuffer format
+ * @param width Framebuffer width
+ * @param height Framebuffer height
+ * @param broadcast The quad id being broadcast
+ */
+function checkFragment(
+data,
+format,
+width,
+height,
+broadcast)
+{
+  assert(broadcast === Math.trunc(broadcast));
+  assert(broadcast >= 0 && broadcast <= 3);
+
+  if (width < 3 || height < 3) {
+    return new Error(
+      `Insufficient framebuffer size [${width}w x ${height}h]. Minimum is [3w x 3h].`
+    );
+  }
+
+  const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format];
+  const blocksPerRow = width / blockWidth;
+  // 256 minimum comes from image copy requirements.
+  const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256);
+  const uintsPerRow = bytesPerRow / 4;
+  const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4;
+
+  const coordToIndex = (row, col) => {
+    return uintsPerRow * row + col * uintsPerTexel;
+  };
+
+  // Iteration skips last row and column to avoid helper invocations because it is not
+  // guaranteed whether or not they participate in the subgroup operation.
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = coordToIndex(row, col);
+
+      const row_is_odd = row % 2 === 1;
+      const col_is_odd = col % 2 === 1;
+
+      // Skip checking quads that extend into potential helper invocations.
+      const max_row = row_is_odd ? row : row + 1;
+      const max_col = col_is_odd ? col : col + 1;
+      if (max_row === height - 1 || max_col === width - 1) {
+        continue;
+      }
+
+      let expect_row = row;
+      let expect_col = col;
+      switch (broadcast) {
+        case 0:
+          expect_row = row_is_odd ? row - 1 : row;
+          expect_col = col_is_odd ? col - 1 : col;
+          break;
+        case 1:
+          expect_row = row_is_odd ? row - 1 : row;
+          expect_col = col_is_odd ? col : col + 1;
+          break;
+        case 2:
+          expect_row = row_is_odd ? row : row + 1;
+          expect_col = col_is_odd ? col - 1 : col;
+          break;
+        case 3:
+          expect_row = row_is_odd ? row : row + 1;
+          expect_col = col_is_odd ? col : col + 1;
+          break;
+      }
+
+      const row_broadcast = data[offset + 1];
+      const col_broadcast = data[offset];
+      if (expect_row !== row_broadcast) {
+        return new Error(`Row ${row}, col ${col}: incorrect row results:
+- expected: ${expect_row}
+-      got: ${row_broadcast}`);
+      }
+
+      if (expect_col !== col_broadcast) {
+        return new Error(`Row ${row}, col ${col}: incorrect col results:
+- expected: ${expect_row}
+-      got: ${col_broadcast}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('fragment,all_active').
+desc(`Tests quadBroadcast in fragment shaders`).
+params((u) =>
+u.
+combine('size', kFramebufferSizes).
+beginSubcases().
+combine('id', [0, 1, 2, 3]).
+combineWithParams([{ format: 'rgba32uint' }])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const fsShader = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> inputs : array<u32>; // unused
+
+@fragment
+fn main(
+  @builtin(position) pos : vec4f,
+) -> @location(0) vec4u {
+  // Force usage
+  _ = inputs[0];
+
+  let linear = u32(pos.x) + u32(pos.y) * ${t.params.size[0]};
+
+  // Filter out possible helper invocations.
+  let x_in_range = u32(pos.x) < (${t.params.size[0]} - 1);
+  let y_in_range = u32(pos.y) < (${t.params.size[1]} - 1);
+  let in_range = x_in_range && y_in_range;
+
+  var x_broadcast = select(1001, u32(pos.x), in_range);
+  var y_broadcast = select(1001, u32(pos.y), in_range);
+
+  x_broadcast = quadBroadcast(x_broadcast, ${t.params.id});
+  y_broadcast = quadBroadcast(y_broadcast, ${t.params.id});
+
+  return vec4u(x_broadcast, y_broadcast, 0, 0);
+}`;
+
+  await runFragmentTest(
+    t,
+    t.params.format,
+    fsShader,
+    t.params.size[0],
+    t.params.size[1],
+    new Uint32Array([0]), // unused,
+    (data) => {
+      return checkFragment(
+        data,
+        t.params.format,
+        t.params.size[0],
+        t.params.size[1],
+        t.params.id
+      );
+    }
+  );
+});
+
+g.test('fragment,split').unimplemented();

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/quadSwap.spec.js

@@ -0,0 +1,666 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for quadSwapX, quadSwapY, and quadSwapDiagnoal.
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
+import { assert, unreachable } from '../../../../../../common/util/util.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import { kBit } from '../../../../../util/constants.js';
+import {
+  kConcreteNumericScalarsAndVectors,
+  Type,
+  VectorType,
+  scalarTypeOf } from
+'../../../../../util/conversion.js';
+import { align } from '../../../../../util/math.js';
+
+import {
+  kWGSizes,
+  kDataSentinel,
+  kPredicateCases,
+  runComputeTest,
+  SubgroupTest,
+  kFramebufferSizes,
+  runFragmentTest } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(SubgroupTest);
+
+const kTypes = objectsToRecord(kConcreteNumericScalarsAndVectors);
+
+
+
+const kOps = ['quadSwapX', 'quadSwapY', 'quadSwapDiagonal'];
+
+/**
+ * Generates scalar values for type
+ *
+ * Generates 4 32-bit values whose bit patterns represent
+ * interesting values of the data type.
+ * @param type The data type
+ */
+function generateScalarValues(type) {
+  const scalarTy = scalarTypeOf(type);
+  switch (scalarTy) {
+    case Type.u32:
+      return [kBit.u32.min, kBit.u32.max, 1111, 2222];
+    case Type.i32:
+      return [
+      kBit.i32.positive.min,
+      kBit.i32.positive.max,
+      kBit.i32.negative.min,
+      0xffffffff // -1
+      ];
+    case Type.f32:
+      return [
+      kBit.f32.positive.zero,
+      kBit.f32.positive.nearest_max,
+      kBit.f32.negative.nearest_min,
+      0xbf800000 // -1
+      ];
+    case Type.f16:
+      return [
+      kBit.f16.positive.zero,
+      kBit.f16.positive.nearest_max,
+      kBit.f16.negative.nearest_min,
+      0xbc00 // -1
+      ];
+    default:
+      unreachable(`Unsupported type: ${type.toString()}`);
+  }
+  return [0, 0, 0, 0];
+}
+
+/**
+ * Generates input bit patterns for the input type
+ *
+ * Generates 4 values of type in a Uint32Array.
+ * 16-bit types are appropriately packed.
+ * @param type The data type
+ */
+function generateTypedInputs(type) {
+  const scalarValues = generateScalarValues(type);
+  let elements = 1;
+  if (type instanceof VectorType) {
+    elements = type.width;
+  }
+  if (type.requiresF16()) {
+    switch (elements) {
+      case 1:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[3] << 16]
+        );
+      case 2:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[3] | scalarValues[3] << 16]
+        );
+      case 3:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[0] | kDataSentinel << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[1] | kDataSentinel << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[2] | kDataSentinel << 16,
+        scalarValues[3] | scalarValues[3] << 16,
+        scalarValues[3] | kDataSentinel << 16]
+        );
+      case 4:
+        return new Uint32Array([
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[0] | scalarValues[0] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[1] | scalarValues[1] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[2] | scalarValues[2] << 16,
+        scalarValues[3] | scalarValues[3] << 16,
+        scalarValues[3] | scalarValues[3] << 16]
+        );
+      default:
+        unreachable(`Unsupported type: ${type.toString()}`);
+    }
+    return new Uint32Array([0]);
+  } else {
+    const bound = elements === 3 ? 4 : elements;
+    const values = [];
+    for (let i = 0; i < 4; i++) {
+      for (let j = 0; j < bound; j++) {
+        if (j < elements) {
+          values.push(scalarValues[i]);
+        } else {
+          values.push(kDataSentinel);
+        }
+      }
+    }
+    return new Uint32Array(values);
+  }
+}
+
+/**
+ * Returns the swapped quad invocation id for the given op
+ *
+ * @param index The index in the range [0,3]
+ * @param op The swap
+ */
+function swapIndex(index, op) {
+  assert(index === Math.trunc(index));
+  assert(index >= 0 && index <= 3);
+  switch (op) {
+    case 'quadSwapX':
+      return index ^ 1;
+    case 'quadSwapY':
+      return index ^ 2;
+    case 'quadSwapDiagonal':
+      return index ^ 3;
+  }
+  unreachable(`Unhandled op ${op}`);
+}
+
+/**
+ * Checks the results of data types test
+ *
+ * The outputs for a given index are expected to match the input values
+ * for the given swap.
+ * @param metadata An unused parameter
+ * @param output The output data
+ * @param input The input data
+ * @param op The type of swap
+ * @param type The data type
+ */
+function checkDataTypes(
+metadata, // unused
+output,
+input,
+op,
+type)
+{
+  if (type.requiresF16() && !(type instanceof VectorType)) {
+    for (let i = 0; i < 4; i++) {
+      const swapIdx = swapIndex(i, op);
+
+      const expectIdx = Math.floor(swapIdx / 2);
+      const expectShift = swapIdx % 2 === 1;
+      let expect = input[expectIdx];
+      if (expectShift) {
+        expect >>= 16;
+      }
+      expect &= 0xffff;
+
+      const resIdx = Math.floor(i / 2);
+      const resShift = i % 2 === 1;
+      let res = output[resIdx];
+      if (resShift) {
+        res >>= 16;
+      }
+      res &= 0xffff;
+
+      if (res !== expect) {
+        return new Error(`${i}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+      }
+    }
+  } else {
+    let uints = 1;
+    if (type instanceof VectorType) {
+      uints = type.width === 3 ? 4 : type.width;
+      if (type.requiresF16()) {
+        uints = Math.floor(uints / 2);
+      }
+    }
+    for (let i = 0; i < 4; i++) {
+      for (let j = 0; j < uints; j++) {
+        const expect = input[swapIndex(i, op) * uints + j];
+        const res = output[i * uints + j];
+        if (res !== expect) {
+          return new Error(`${uints * i + j}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+        }
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('data_types').
+desc('Test allowed data types').
+params((u) => u.combine('type', keysOf(kTypes)).beginSubcases().combine('op', kOps)).
+beforeAllSubcases((t) => {
+  const features = ['subgroups'];
+  const type = kTypes[t.params.type];
+  if (type.requiresF16()) {
+    features.push('subgroups-f16');
+    features.push('shader-f16');
+  }
+  t.selectDeviceOrSkipTestCase(features);
+}).
+fn(async (t) => {
+  const wgSize = [4, 1, 1];
+  const type = kTypes[t.params.type];
+  let enables = `enable subgroups;\n`;
+  if (type.requiresF16()) {
+    enables += `enable f16;\nenable subgroups_f16;`;
+  }
+  const wgsl = `
+${enables}
+
+@group(0) @binding(0)
+var<storage> input : array<${type.toString()}>;
+
+@group(0) @binding(1)
+var<storage, read_write> output : array<${type.toString()}>;
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : array<u32>; // unused
+
+@compute @workgroup_size(${wgSize[0]}, ${wgSize[1]}, ${wgSize[2]})
+fn main(
+  @builtin(subgroup_invocation_id) id : u32,
+) {
+  // Force usage
+  _ = metadata[0];
+
+  output[id] = ${t.params.op}(input[id]);
+}`;
+
+  const inputData = generateTypedInputs(type);
+  let uintsPerOutput = 1;
+  if (type instanceof VectorType) {
+    uintsPerOutput = type.width === 3 ? 4 : type.width;
+    if (type.requiresF16()) {
+      uintsPerOutput = Math.floor(uintsPerOutput / 2);
+    }
+  }
+  await runComputeTest(
+    t,
+    wgsl,
+    wgSize,
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkDataTypes(metadata, output, inputData, t.params.op, type);
+    }
+  );
+});
+
+/**
+ * Checks quad swaps in compute shaders
+ *
+ * Assumes that quads are linear within a subgroup.
+ *
+ * @param metadata An array of integers divided as follows:
+ *                 * first half subgroup invocation ids
+ *                 * second half subgroup sizes
+ * @param output An array of integers divided as follows:
+ *               * first half results of quad broadcast
+ *               * second half generated unique subgroup ids
+ * @param op The swap operation
+ * @param filter A functor to filter active invocations
+ */
+function checkSwapCompute(
+metadata,
+output,
+op,
+filter)
+{
+  const bound = Math.floor(output.length / 2);
+  for (let i = 0; i < bound; i++) {
+    const subgroup_id = output[bound + i];
+    const id = metadata[i];
+    const size = metadata[bound + i];
+    if (!filter(id, size)) {
+      if (output[i] !== kDataSentinel) {
+        return new Error(`Unexpected write for invocation ${i}`);
+      }
+      continue;
+    }
+
+    const quad_id = Math.floor(id / 4);
+    const quad_index = id % 4;
+    let found = false;
+    for (let j = 0; j < bound; j++) {
+      const other_id = metadata[j];
+      const other_quad_id = Math.floor(other_id / 4);
+      const other_quad_index = other_id % 4;
+      const other_subgroup_id = output[bound + j];
+      if (
+      subgroup_id === other_subgroup_id &&
+      quad_id === other_quad_id &&
+      quad_index === swapIndex(other_quad_index, op))
+      {
+        found = true;
+        if (output[i] !== j) {
+          return new Error(`Invocation ${i}: incorrect result
+- expected: ${j}
+-      got: ${output[i]}`);
+        }
+        break;
+      }
+    }
+    if (!found) {
+      return new Error(`Invocation ${i}: failed to find swapped result`);
+    }
+  }
+
+  return undefined;
+}
+
+g.test('compute,all_active').
+desc(
+  `Tests swaps with all active invocations
+
+Quad operations require a full quad so workgroup sizes are limited to multiples of 4.
+  `
+).
+params((u) =>
+u.
+combine('wgSize', kWGSizes).
+filter((t) => {
+  const wgThreads = t.wgSize[0] * t.wgSize[1] * t.wgSize[2];
+  return wgThreads % 4 === 0;
+}).
+beginSubcases().
+combine('op', kOps)
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : u32; // unused
+
+struct Output {
+  results : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(1)
+var<storage, read_write> output : Output;
+
+struct Metadata {
+  id : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  // Force usage
+  _ = inputs;
+
+  let b = ${t.params.op}(lid);
+  output.results[lid] = b;
+  output.subgroup_size[lid] = subgroupBroadcastFirst(lid + 1);
+  metadata.id[lid] = id;
+  metadata.subgroup_size[lid] = subgroupSize;
+}`;
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    new Uint32Array([0]), // unused
+    (metadata, output) => {
+      return checkSwapCompute(metadata, output, t.params.op, (id, size) => {
+        return true;
+      });
+    }
+  );
+});
+
+g.test('compute,split').
+desc(
+  `Tests swaps with all predicated invocations
+
+Quad operations require a full quad so workgroup sizes are limited to multiples of 4.
+Quad operations require a fully active quad to operate correctly so several of the
+predication filters are skipped.
+  `
+).
+params((u) =>
+u.
+combine('predicate', keysOf(kPredicateCases)).
+filter((t) => {
+  return t.predicate === 'lower_half' || t.predicate === 'upper_half';
+}).
+combine('wgSize', kWGSizes).
+filter((t) => {
+  const wgThreads = t.wgSize[0] * t.wgSize[1] * t.wgSize[2];
+  return wgThreads % 4 === 0;
+}).
+beginSubcases().
+combine('op', kOps)
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+  const testcase = kPredicateCases[t.params.predicate];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : u32; // unused
+
+struct Output {
+  results : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(1)
+var<storage, read_write> output : Output;
+
+struct Metadata {
+  id : array<u32, ${wgThreads}>,
+  subgroup_size : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  // Force usage
+  _ = inputs;
+
+  output.subgroup_size[lid] = subgroupBroadcastFirst(lid + 1);
+  metadata.id[lid] = id;
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  if ${testcase.cond} {
+    let b = ${t.params.op}(lid);
+    output.results[lid] = b;
+  }
+}`;
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    new Uint32Array([0]), // unused
+    (metadata, output) => {
+      return checkSwapCompute(metadata, output, t.params.op, testcase.filter);
+    }
+  );
+});
+
+/**
+ * Checks results of quad swaps in fragment shaders.
+ *
+ * @param data The framebuffer output
+ *             * component 0 is the broadcast of the integer x position
+ *             * component 1 is the broadcast of the integer y position
+ * @param format The framebuffer format
+ * @param width Framebuffer width
+ * @param height Framebuffer height
+ * @param broadcast The quad id being broadcast
+ */
+function checkFragment(
+data,
+format,
+width,
+height,
+op)
+{
+  if (width < 3 || height < 3) {
+    return new Error(
+      `Insufficient framebuffer size [${width}w x ${height}h]. Minimum is [3w x 3h].`
+    );
+  }
+
+  const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format];
+  const blocksPerRow = width / blockWidth;
+  // 256 minimum comes from image copy requirements.
+  const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256);
+  const uintsPerRow = bytesPerRow / 4;
+  const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4;
+
+  const coordToIndex = (row, col) => {
+    return uintsPerRow * row + col * uintsPerTexel;
+  };
+
+  // Iteration skips last row and column to avoid helper invocations because it is not
+  // guaranteed whether or not they participate in the subgroup operation.
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = coordToIndex(row, col);
+
+      const row_is_odd = row % 2 === 1;
+      const col_is_odd = col % 2 === 1;
+
+      // Skip checking quads that extend into potential helper invocations.
+      const max_row = row_is_odd ? row : row + 1;
+      const max_col = col_is_odd ? col : col + 1;
+      if (max_row === height - 1 || max_col === width - 1) {
+        continue;
+      }
+
+      let expect_row = row;
+      let expect_col = col;
+      switch (op) {
+        case 'quadSwapX':
+          expect_col = col_is_odd ? col - 1 : col + 1;
+          break;
+        case 'quadSwapY':
+          expect_row = row_is_odd ? row - 1 : row + 1;
+          break;
+        case 'quadSwapDiagonal':
+          expect_row = row_is_odd ? row - 1 : row + 1;
+          expect_col = col_is_odd ? col - 1 : col + 1;
+          break;
+      }
+
+      const row_output = data[offset + 1];
+      const col_output = data[offset];
+      if (expect_row !== row_output) {
+        return new Error(`Row ${row}, col ${col}: incorrect row results:
+- expected: ${expect_row}
+-      got: ${row_output}`);
+      }
+
+      if (expect_col !== col_output) {
+        return new Error(`Row ${row}, col ${col}: incorrect col results:
+- expected: ${expect_row}
+-      got: ${col_output}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('fragment,all_active').
+desc(`Tests quad swaps in fragment shaders`).
+params((u) =>
+u.
+combine('size', kFramebufferSizes).
+beginSubcases().
+combine('op', kOps).
+combineWithParams([{ format: 'rgba32uint' }])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const fsShader = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> inputs : array<u32>; // unused
+
+@fragment
+fn main(
+  @builtin(position) pos : vec4f,
+) -> @location(0) vec4u {
+  // Force usage
+  _ = inputs[0];
+
+  let linear = u32(pos.x) + u32(pos.y) * ${t.params.size[0]};
+
+  // Filter out possible helper invocations.
+  let x_in_range = u32(pos.x) < (${t.params.size[0]} - 1);
+  let y_in_range = u32(pos.y) < (${t.params.size[1]} - 1);
+  let in_range = x_in_range && y_in_range;
+
+  var x_swap = select(1001, u32(pos.x), in_range);
+  var y_swap = select(1001, u32(pos.y), in_range);
+
+  x_swap = ${t.params.op}(x_swap);
+  y_swap = ${t.params.op}(y_swap);
+
+  return vec4u(x_swap, y_swap, 0, 0);
+}`;
+
+  await runFragmentTest(
+    t,
+    t.params.format,
+    fsShader,
+    t.params.size[0],
+    t.params.size[1],
+    new Uint32Array([0]), // unused,
+    (data) => {
+      return checkFragment(
+        data,
+        t.params.format,
+        t.params.size[0],
+        t.params.size[1],
+        t.params.op
+      );
+    }
+  );
+});
+
+g.test('fragment,split').unimplemented();

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/smoothstep.spec.js

@@ -9,9 +9,14 @@ T is S or vecN<S>
 Returns the smooth Hermite interpolation between 0 and 1.
 Component-wise when T is a vector.
 For scalar T, the result is t * t * (3.0 - 2.0 * t), where t = clamp((x - low) / (high - low), 0.0, 1.0).
+
+If low >= high:
+* It is a shader-creation error if low and high are const-expressions.
+* It is a pipeline-creation error if low and high are override-expressions.
 `;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
 import { GPUTest } from '../../../../../gpu_test.js';
 import { Type } from '../../../../../util/conversion.js';
+
 import { allInputSources, onlyConstInputSource, run } from '../../expression.js';
 
 import { abstractFloatBuiltin, builtin } from './builtin.js';
@@ -19,6 +24,13 @@ import { d } from './smoothstep.cache.js';
 
 export const g = makeTestGroup(GPUTest);
 
+// Returns true if `c` is valid for a const evaluation of smoothstep.
+function validForConst(c) {
+  const low = c.input[0];
+  const high = c.input[1];
+  return low.value < high.value;
+}
+
 g.test('abstract_float').
 specURL('https://www.w3.org/TR/WGSL/#float-builtin-functions').
 desc(`abstract float tests`).
@@ -28,7 +40,7 @@ combine('inputSource', onlyConstInputSource).
 combine('vectorize', [undefined, 2, 3, 4])
 ).
 fn(async (t) => {
-  const cases = await d.get('abstract_const');
+  const cases = (await d.get('abstract_const')).filter((c) => validForConst(c));
   await run(
     t,
     abstractFloatBuiltin('smoothstep'),
@@ -47,7 +59,15 @@ u.combine('inputSource', allInputSources).combine('vectorize', [undefined, 2, 3,
 ).
 fn(async (t) => {
   const cases = await d.get(t.params.inputSource === 'const' ? 'f32_const' : 'f32_non_const');
-  await run(t, builtin('smoothstep'), [Type.f32, Type.f32, Type.f32], Type.f32, t.params, cases);
+  const validCases = cases.filter((c) => t.params.inputSource !== 'const' || validForConst(c));
+  await run(
+    t,
+    builtin('smoothstep'),
+    [Type.f32, Type.f32, Type.f32],
+    Type.f32,
+    t.params,
+    validCases
+  );
 });
 
 g.test('f16').
@@ -61,5 +81,13 @@ beforeAllSubcases((t) => {
 }).
 fn(async (t) => {
   const cases = await d.get(t.params.inputSource === 'const' ? 'f16_const' : 'f16_non_const');
-  await run(t, builtin('smoothstep'), [Type.f16, Type.f16, Type.f16], Type.f16, t.params, cases);
+  const validCases = cases.filter((c) => t.params.inputSource !== 'const' || validForConst(c));
+  await run(
+    t,
+    builtin('smoothstep'),
+    [Type.f16, Type.f16, Type.f16],
+    Type.f16,
+    t.params,
+    validCases
+  );
 });

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/subgroupAdd.spec.js

@@ -0,0 +1,364 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for subgroupAdd, subgroupExclusiveAdd, and subgroupInclusiveAdd
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
+import { iterRange } from '../../../../../../common/util/util.js';
+import { GPUTest } from '../../../../../gpu_test.js';
+import {
+  kConcreteNumericScalarsAndVectors,
+  Type,
+  VectorType,
+  numberToFloatBits,
+  floatBitsToNumber,
+  kFloat32Format,
+  kFloat16Format,
+  scalarTypeOf } from
+'../../../../../util/conversion.js';
+import { FP } from '../../../../../util/floating_point.js';
+
+import {
+  kNumCases,
+  kStride,
+  kWGSizes,
+  kPredicateCases,
+  runAccuracyTest,
+  runComputeTest } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(GPUTest);
+
+const kIdentity = 0;
+
+const kDataTypes = objectsToRecord(kConcreteNumericScalarsAndVectors);
+
+const kOperations = ['subgroupAdd', 'subgroupExclusiveAdd', 'subgroupInclusiveAdd'];
+
+g.test('fp_accuracy').
+desc(
+  `Tests the accuracy of floating-point addition.
+
+The order of operations is implementation defined, most threads are filled with
+the identity value and two receive random values.
+Subgroup sizes are not known ahead of time so some cases may not perform any
+interesting operations. The test biases towards checking subgroup sizes under 64.
+These tests only check two values in order to reuse more of the existing infrastructure
+and limit the number of permutations needed to calculate the final result.`
+).
+params((u) =>
+u.
+combine('case', [...iterRange(kNumCases, (x) => x)]).
+combine('type', ['f32', 'f16']).
+combine('wgSize', [
+[kStride, 1, 1],
+[kStride / 2, 2, 1]]
+)
+).
+beforeAllSubcases((t) => {
+  const features = ['subgroups'];
+  if (t.params.type === 'f16') {
+    features.push('shader-f16');
+    features.push('subgroups-f16');
+  }
+  t.selectDeviceOrSkipTestCase(features);
+}).
+fn(async (t) => {
+  await runAccuracyTest(
+    t,
+    t.params.case,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    'subgroupAdd',
+    t.params.type,
+    kIdentity,
+    t.params.type === 'f16' ? FP.f16.additionInterval : FP.f32.additionInterval
+  );
+});
+
+/**
+ * Checks subgroup additions
+ *
+ * Expected results:
+ * - subgroupAdd: each invocation should have result equal to real subgroup size
+ * - subgroupExclusiveAdd: each invocation should have result equal to its subgroup invocation id
+ * - subgroupInclusiveAdd: each invocation should be equal to the result of subgroupExclusiveAdd plus the fill value
+ * @param metadata An array containing actual subgroup size per invocation followed by
+ *                 subgroup invocation id per invocation
+ * @param output An array of additions
+ * @param type The data type
+ * @param operation Type of addition
+ * @param expectedfillValue The original value used to fill the test array
+ */
+function checkAddition(
+metadata,
+output,
+type,
+operation,
+expectedfillValue)
+{
+  let numEles = 1;
+  if (type instanceof VectorType) {
+    numEles = type.width;
+  }
+  const scalarTy = scalarTypeOf(type);
+  const expectedOffset = operation === 'subgroupAdd' ? 0 : metadata.length / 2;
+  for (let i = 0; i < metadata.length / 2; i++) {
+    let expected = metadata[i + expectedOffset];
+    if (operation === 'subgroupInclusiveAdd') {
+      expected += expectedfillValue;
+    }
+
+    for (let j = 0; j < numEles; j++) {
+      let idx = i * numEles + j;
+      const isOdd = idx & 0x1;
+      if (scalarTy === Type.f16) {
+        idx = Math.floor(idx / 2);
+      }
+      let val = output[idx];
+      if (scalarTy === Type.f32) {
+        val = floatBitsToNumber(val, kFloat32Format);
+      } else if (scalarTy === Type.f16) {
+        if (isOdd) {
+          val = val >> 16;
+        }
+        val = floatBitsToNumber(val & 0xffff, kFloat16Format);
+      }
+      if (expected !== val) {
+        return new Error(`Invocation ${i}, component ${j}: incorrect result
+- expected: ${expected}
+-      got: ${val}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('data_types').
+desc(
+  `Tests subgroup addition for valid data types
+
+Tests a simple addition of all 1 values.
+Reductions expect result to be equal to actual subgroup size.
+Exclusice scans expect result to be equal subgroup invocation id.
+
+TODO: support vec3 types.
+  `
+).
+params((u) =>
+u.
+combine('type', keysOf(kDataTypes)).
+filter((t) => {
+  const type = kDataTypes[t.type];
+  if (type instanceof VectorType) {
+    return type.width !== 3;
+  }
+  return true;
+}).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('operation', kOperations)
+).
+beforeAllSubcases((t) => {
+  const features = ['subgroups'];
+  const type = kDataTypes[t.params.type];
+  if (type.requiresF16()) {
+    features.push('shader-f16');
+    features.push('subgroups-f16');
+  }
+  t.selectDeviceOrSkipTestCase(features);
+}).
+fn(async (t) => {
+  const type = kDataTypes[t.params.type];
+  let numEles = 1;
+  if (type instanceof VectorType) {
+    numEles = type.width;
+  }
+  const scalarType = scalarTypeOf(type);
+  let enables = 'enable subgroups;\n';
+  if (type.requiresF16()) {
+    enables += 'enable f16;\nenable subgroups_f16;\n';
+  }
+
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+${enables}
+
+@group(0) @binding(0)
+var<storage> inputs : array<${type.toString()}>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<${type.toString()}>;
+
+struct Metadata {
+  subgroup_size : array<u32, ${wgThreads}>,
+  subgroup_invocation_id : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+) {
+  // Record the actual subgroup size for this invocation.
+  // Note: subgroup_size builtin is always a power-of-2 and might be larger
+  // if the subgroup is not full.
+  let ballot = subgroupBallot(true);
+  var size = countOneBits(ballot.x);
+  size += countOneBits(ballot.y);
+  size += countOneBits(ballot.z);
+  size += countOneBits(ballot.w);
+  metadata.subgroup_size[lid] = size;
+
+  // Record subgroup invocation id for this invocation.
+  metadata.subgroup_invocation_id[lid] = id;
+
+  outputs[lid] = ${t.params.operation}(inputs[lid]);
+}`;
+  const expectedFillValue = 1;
+  let fillValue = expectedFillValue;
+  let numUints = wgThreads * numEles;
+  if (scalarType === Type.f32) {
+    fillValue = numberToFloatBits(1, kFloat32Format);
+  } else if (scalarType === Type.f16) {
+    const f16 = numberToFloatBits(1, kFloat16Format);
+    fillValue = f16 | f16 << 16;
+    numUints = Math.ceil(numUints / 2);
+  }
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    numUints,
+    new Uint32Array([...iterRange(numUints, (x) => fillValue)]),
+    (metadata, output) => {
+      return checkAddition(metadata, output, type, t.params.operation, expectedFillValue);
+    }
+  );
+});
+
+g.test('fragment').unimplemented();
+
+/**
+ * Performs correctness checking for predicated additions
+ *
+ * Assumes the shader performs a predicated subgroup addition with the
+ * subgroup_invocation_id as the data.
+ *
+ * @param metadata An array containing subgroup sizes and subgroup invocation ids
+ * @param output An array containing the output results
+ * @param operation The type of addition
+ * @param filter A functor that mirrors the predication in the shader
+ */
+function checkPredicatedAddition(
+metadata,
+output,
+operation,
+filter)
+{
+  for (let i = 0; i < output.length; i++) {
+    const size = metadata[i];
+    const id = metadata[output.length + i];
+    let expected = 0;
+    if (filter(id, size)) {
+      const bound =
+      operation === 'subgroupInclusiveAdd' ? id + 1 : operation === 'subgroupAdd' ? size : id;
+      for (let j = 0; j < bound; j++) {
+        if (filter(j, size)) {
+          expected += j;
+        }
+      }
+    } else {
+      expected = 999;
+    }
+    if (expected !== output[i]) {
+      return new Error(`Invocation ${i}: incorrect result
+- expected: ${expected}
+-      got: ${output[i]}`);
+    }
+  }
+  return undefined;
+}
+
+g.test('compute,split').
+desc('Tests that only active invocations contribute to the operation').
+params((u) =>
+u.
+combine('case', keysOf(kPredicateCases)).
+beginSubcases().
+combine('operation', kOperations).
+combine('wgSize', kWGSizes)
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kPredicateCases[t.params.case];
+  const outputUintsPerElement = 1;
+  const inputData = new Uint32Array([0]); // no input data
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> input : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  subgroup_size : array<u32, ${wgThreads}>,
+  subgroup_invocation_id : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+) {
+  _ = input[0];
+
+  // Record the actual subgroup size for this invocation.
+  // Note: subgroup_size builtin is always a power-of-2 and might be larger
+  // if the subgroup is not full.
+  let ballot = subgroupBallot(true);
+  var subgroupSize = countOneBits(ballot.x);
+  subgroupSize += countOneBits(ballot.y);
+  subgroupSize += countOneBits(ballot.z);
+  subgroupSize += countOneBits(ballot.w);
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  // Record subgroup invocation id for this invocation.
+  metadata.subgroup_invocation_id[lid] = id;
+
+  if ${testcase.cond} {
+    outputs[lid] = ${t.params.operation}(id);
+  } else {
+    return;
+  }
+}`;
+
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    outputUintsPerElement,
+    inputData,
+    (metadata, output) => {
+      return checkPredicatedAddition(metadata, output, t.params.operation, testcase.filter);
+    }
+  );
+});

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/subgroupAll.spec.js

@@ -0,0 +1,390 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for subgroupAll.
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf } from '../../../../../../common/util/data_tables.js';
+import { iterRange } from '../../../../../../common/util/util.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import { align } from '../../../../../util/math.js';
+import { PRNG } from '../../../../../util/prng.js';
+
+import {
+  kWGSizes,
+  kPredicateCases,
+  SubgroupTest,
+  kDataSentinel,
+  kFramebufferSizes,
+  runComputeTest,
+  runFragmentTest } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(SubgroupTest);
+
+const kNumCases = 15;
+
+/**
+ * Generate input data for testing.
+ *
+ * Data is generated in the following categories:
+ * Seed 0 generates all 0 data
+ * Seed 1 generates all 1 data
+ * Seeds 2-9 generates all 1s except for a zero randomly once per 32 elements
+ * Seeds 10+ generate all random data
+ * @param seed The seed for the PRNG
+ * @param num The number of data items to generate
+ */
+function generateInputData(seed, num) {
+  const prng = new PRNG(seed);
+
+  const bound = Math.min(num, 32);
+  const index = prng.uniformInt(bound);
+
+  return new Uint32Array([
+  ...iterRange(num, (x) => {
+    if (seed === 0) {
+      return 0;
+    } else if (seed === 1) {
+      return 1;
+    } else if (seed < 10) {
+      const bounded = x % bound;
+      return bounded === index ? 0 : 1;
+    }
+    return prng.uniformInt(2);
+  })]
+  );
+}
+
+/**
+ * Checks the result of a subgroupAll operation
+ *
+ * Since subgroup size depends on the pipeline compile, we calculate the expected
+ * results after execution. The shader generates a subgroup id and records it for
+ * each invocation. The check first calculates the expected result for each subgroup
+ * and then compares to the actual result for each invocation. The filter functor
+ * ensures only the correct invocations contribute to the calculation.
+ * @param metadata An array of uints:
+ *                 * first half containing subgroup sizes (from builtin value)
+ *                 * second half subgroup invocation id
+ * @param output An array of uints containing:
+ *               * first half is the outputs of subgroupAll
+ *               * second half is a generated subgroup id
+ * @param numInvs Number of invocations executed
+ * @param input The input data (equal size to output)
+ * @param filter A functor to filter active invocations
+ */
+function checkAll(
+metadata, // unused
+output,
+numInvs,
+input,
+filter)
+{
+  // First, generate expected results.
+  const expected = new Map();
+  for (let inv = 0; inv < numInvs; inv++) {
+    const size = metadata[inv];
+    const id = metadata[inv + numInvs];
+    if (!filter(id, size)) {
+      continue;
+    }
+    const subgroup_id = output[numInvs + inv];
+    let v = expected.get(subgroup_id) ?? 1;
+    v &= input[inv];
+    expected.set(subgroup_id, v);
+  }
+
+  // Second, check against actual results.
+  for (let inv = 0; inv < numInvs; inv++) {
+    const size = metadata[inv];
+    const id = metadata[inv + numInvs];
+    const res = output[inv];
+    if (filter(id, size)) {
+      const subgroup_id = output[numInvs + inv];
+      const expected_v = expected.get(subgroup_id) ?? 0;
+      if (expected_v !== res) {
+        return new Error(`Invocation ${inv}:
+- expected: ${expected_v}
+-      got: ${res}`);
+      }
+    } else {
+      if (res !== kDataSentinel) {
+        return new Error(`Invocation ${inv} unexpected write:
+- subgroup invocation id: ${id}
+-          subgroup size: ${size}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('compute,all_active').
+desc(`Test compute subgroupAll`).
+params((u) =>
+u.
+combine('wgSize', kWGSizes).
+beginSubcases().
+combine('case', [...iterRange(kNumCases, (x) => x)])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  subgroup_size: array<u32, ${wgThreads}>,
+  subgroup_invocation_id: array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  metadata.subgroup_invocation_id[lid] = id;
+
+  // Record a representative subgroup id.
+  outputs[lid + ${wgThreads}] = subgroupBroadcastFirst(lid);
+
+  let res = select(0u, 1u, subgroupAll(bool(inputs[lid])));
+  outputs[lid] = res;
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads);
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkAll(metadata, output, wgThreads, inputData, (id, size) => {
+        return true;
+      });
+    }
+  );
+});
+
+g.test('compute,split').
+desc('Test that only active invocation participate').
+params((u) =>
+u.
+combine('predicate', keysOf(kPredicateCases)).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('case', [...iterRange(kNumCases, (x) => x)])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kPredicateCases[t.params.predicate];
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  subgroup_size : array<u32, ${wgThreads}>,
+  subgroup_invocation_id : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  // Record subgroup invocation id for this invocation.
+  metadata.subgroup_invocation_id[lid] = id;
+
+  // Record a generated subgroup id.
+  outputs[${wgThreads} + lid] = subgroupBroadcastFirst(lid);
+
+  if ${testcase.cond} {
+    outputs[lid] = select(0u, 1u, subgroupAll(bool(inputs[lid])));
+  } else {
+    return;
+  }
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads);
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkAll(metadata, output, wgThreads, inputData, testcase.filter);
+    }
+  );
+});
+
+/**
+ * Checks subgroupAll results from a fragment shader.
+ *
+ * @param data Framebuffer output
+ *             * component 0 is result
+ *             * component 1 is generated subgroup id
+ * @param input An array of input data
+ * @param format The framebuffer format
+ * @param width Framebuffer width
+ * @param height Framebuffer height
+ */
+function checkFragmentAll(
+data,
+input,
+format,
+width,
+height)
+{
+  const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format];
+  const blocksPerRow = width / blockWidth;
+  // 256 minimum comes from image copy requirements.
+  const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256);
+  const uintsPerRow = bytesPerRow / 4;
+  const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4;
+
+  // Iteration skips last row and column to avoid helper invocations because it is not
+  // guaranteed whether or not they participate in the subgroup operation.
+  const expected = new Map();
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        return new Error(`Internal error: helper invocation at (${col}, ${row})`);
+      }
+
+      let v = expected.get(subgroup_id) ?? 1;
+      // First index of input is an atomic counter.
+      v &= input[row * width + col];
+      expected.set(subgroup_id, v);
+    }
+  }
+
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const res = data[offset];
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        // Inactive in the fragment.
+        continue;
+      }
+
+      const expected_v = expected.get(subgroup_id) ?? 0;
+      if (expected_v !== res) {
+        return new Error(`Row ${row}, col ${col}: incorrect results:
+- expected: ${expected_v}
+-      got: ${res}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('fragment,all_active').
+desc('Tests subgroupAll in fragment shaders').
+params((u) =>
+u.
+combine('size', kFramebufferSizes).
+beginSubcases().
+combine('case', [...iterRange(kNumCases, (x) => x)]).
+combineWithParams([{ format: 'rg32uint' }])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const numInputs = t.params.size[0] * t.params.size[1];
+  const inputData = generateInputData(t.params.case, numInputs);
+
+  const fsShader = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> inputs : array<u32>;
+
+@fragment
+fn main(
+  @builtin(position) pos : vec4f,
+) -> @location(0) vec2u {
+  // Generate a subgroup id based on linearized position, but avoid 0.
+  let linear = u32(pos.x) + u32(pos.y) * ${t.params.size[0]};
+  var subgroup_id = linear + 1;
+  subgroup_id = subgroupBroadcastFirst(subgroup_id);
+
+  // Filter out possible helper invocations.
+  let x_in_range = u32(pos.x) < (${t.params.size[0]} - 1);
+  let y_in_range = u32(pos.y) < (${t.params.size[1]} - 1);
+  let in_range = x_in_range && y_in_range;
+  let input = select(1u, inputs[linear], in_range);
+
+  let res = select(0u, 1u, subgroupAll(bool(input)));
+  return vec2u(res, subgroup_id);
+}`;
+
+  await runFragmentTest(
+    t,
+    t.params.format,
+    fsShader,
+    t.params.size[0],
+    t.params.size[1],
+    inputData,
+    (data) => {
+      return checkFragmentAll(
+        data,
+        inputData,
+        t.params.format,
+        t.params.size[0],
+        t.params.size[1]
+      );
+    }
+  );
+});
+
+// Using subgroup operations in control with fragment shaders
+// quickly leads to unportable behavior.
+g.test('fragment,split').unimplemented();

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/subgroupAny.spec.js

@@ -0,0 +1,390 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for subgroupAny.
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf } from '../../../../../../common/util/data_tables.js';
+import { iterRange } from '../../../../../../common/util/util.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import { align } from '../../../../../util/math.js';
+import { PRNG } from '../../../../../util/prng.js';
+
+import {
+  kWGSizes,
+  kPredicateCases,
+  SubgroupTest,
+  kDataSentinel,
+  runComputeTest,
+  runFragmentTest,
+  kFramebufferSizes } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(SubgroupTest);
+
+const kNumCases = 15;
+
+/**
+ * Generate input data for testing.
+ *
+ * Data is generated in the following categories:
+ * Seed 0 generates all 0 data
+ * Seed 1 generates all 1 data
+ * Seeds 2-9 generates all 0s except for a one randomly once per 32 elements
+ * Seeds 10+ generate all random data
+ * @param seed The seed for the PRNG
+ * @param num The number of data items to generate
+ */
+function generateInputData(seed, num) {
+  const prng = new PRNG(seed);
+
+  const bound = Math.min(num, 32);
+  const index = prng.uniformInt(bound);
+
+  return new Uint32Array([
+  ...iterRange(num, (x) => {
+    if (seed === 0) {
+      return 0;
+    } else if (seed === 1) {
+      return 1;
+    } else if (seed < 10) {
+      const bounded = x % bound;
+      return bounded === index ? 1 : 0;
+    }
+    return prng.uniformInt(2);
+  })]
+  );
+}
+
+/**
+ * Checks the result of a subgroupAny operation
+ *
+ * Since subgroup size depends on the pipeline compile, we calculate the expected
+ * results after execution. The shader generates a subgroup id and records it for
+ * each invocation. The check first calculates the expected result for each subgroup
+ * and then compares to the actual result for each invocation. The filter functor
+ * ensures only the correct invocations contribute to the calculation.
+ * @param metadata An array of uints:
+ *                 * first half containing subgroup sizes (from builtin value)
+ *                 * second half subgroup invocation id
+ * @param output An array of uints containing:
+ *               * first half is the outputs of subgroupAny
+ *               * second half is a generated subgroup id
+ * @param numInvs Number of invocations executed
+ * @param input The input data (equal size to output)
+ * @param filter A functor to filter active invocations
+ */
+function checkAny(
+metadata, // unused
+output,
+numInvs,
+input,
+filter)
+{
+  // First, generate expected results.
+  const expected = new Map();
+  for (let inv = 0; inv < numInvs; inv++) {
+    const size = metadata[inv];
+    const id = metadata[inv + numInvs];
+    if (!filter(id, size)) {
+      continue;
+    }
+    const subgroup_id = output[numInvs + inv];
+    let v = expected.get(subgroup_id) ?? 0;
+    v |= input[inv];
+    expected.set(subgroup_id, v);
+  }
+
+  // Second, check against actual results.
+  for (let inv = 0; inv < numInvs; inv++) {
+    const size = metadata[inv];
+    const id = metadata[inv + numInvs];
+    const res = output[inv];
+    if (filter(id, size)) {
+      const subgroup_id = output[numInvs + inv];
+      const expected_v = expected.get(subgroup_id) ?? 0;
+      if (expected_v !== res) {
+        return new Error(`Invocation ${inv}:
+- expected: ${expected_v}
+-      got: ${res}`);
+      }
+    } else {
+      if (res !== kDataSentinel) {
+        return new Error(`Invocation ${inv} unexpected write:
+- subgroup invocation id: ${id}
+-          subgroup size: ${size}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('compute,all_active').
+desc(`Test compute subgroupAny`).
+params((u) =>
+u.
+combine('wgSize', kWGSizes).
+beginSubcases().
+combine('case', [...iterRange(kNumCases, (x) => x)])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  subgroup_size: array<u32, ${wgThreads}>,
+  subgroup_invocation_id: array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  metadata.subgroup_invocation_id[lid] = id;
+
+  // Record a representative subgroup id.
+  outputs[lid + ${wgThreads}] = subgroupBroadcastFirst(lid);
+
+  let res = select(0u, 1u, subgroupAny(bool(inputs[lid])));
+  outputs[lid] = res;
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads);
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkAny(metadata, output, wgThreads, inputData, (id, size) => {
+        return true;
+      });
+    }
+  );
+});
+
+g.test('compute,split').
+desc('Test that only active invocation participate').
+params((u) =>
+u.
+combine('predicate', keysOf(kPredicateCases)).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('case', [...iterRange(kNumCases, (x) => x)])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kPredicateCases[t.params.predicate];
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  subgroup_size : array<u32, ${wgThreads}>,
+  subgroup_invocation_id : array<u32, ${wgThreads}>,
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+  metadata.subgroup_size[lid] = subgroupSize;
+
+  // Record subgroup invocation id for this invocation.
+  metadata.subgroup_invocation_id[lid] = id;
+
+  // Record a generated subgroup id.
+  outputs[${wgThreads} + lid] = subgroupBroadcastFirst(lid);
+
+  if ${testcase.cond} {
+    outputs[lid] = select(0u, 1u, subgroupAny(bool(inputs[lid])));
+  } else {
+    return;
+  }
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads);
+
+  const uintsPerOutput = 2;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkAny(metadata, output, wgThreads, inputData, testcase.filter);
+    }
+  );
+});
+
+/**
+ * Checks subgroupAny results from a fragment shader.
+ *
+ * @param data Framebuffer output
+ *             * component 0 is result
+ *             * component 1 is generated subgroup id
+ * @param input An array of input data
+ * @param format The framebuffer format
+ * @param width Framebuffer width
+ * @param height Framebuffer height
+ */
+function checkFragmentAny(
+data,
+input,
+format,
+width,
+height)
+{
+  const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format];
+  const blocksPerRow = width / blockWidth;
+  // 256 minimum comes from image copy requirements.
+  const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256);
+  const uintsPerRow = bytesPerRow / 4;
+  const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4;
+
+  // Iteration skips last row and column to avoid helper invocations because it is not
+  // guaranteed whether or not they participate in the subgroup operation.
+  const expected = new Map();
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        return new Error(`Internal error: helper invocation at (${col}, ${row})`);
+      }
+
+      let v = expected.get(subgroup_id) ?? 0;
+      // First index of input is an atomic counter.
+      v |= input[row * width + col];
+      expected.set(subgroup_id, v);
+    }
+  }
+
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const res = data[offset];
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        // Inactive in the fragment.
+        continue;
+      }
+
+      const expected_v = expected.get(subgroup_id) ?? 0;
+      if (expected_v !== res) {
+        return new Error(`Row ${row}, col ${col}: incorrect results:
+- expected: ${expected_v}
+-      got: ${res}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('fragment,all_active').
+desc('Tests subgroupAny in fragment shaders').
+params((u) =>
+u.
+combine('size', kFramebufferSizes).
+beginSubcases().
+combine('case', [...iterRange(kNumCases, (x) => x)]).
+combineWithParams([{ format: 'rg32uint' }])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const numInputs = t.params.size[0] * t.params.size[1];
+  const inputData = generateInputData(t.params.case, numInputs);
+
+  const fsShader = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> inputs : array<u32>;
+
+@fragment
+fn main(
+  @builtin(position) pos : vec4f,
+) -> @location(0) vec2u {
+  // Generate a subgroup id based on linearized position, but avoid 0.
+  let linear = u32(pos.x) + u32(pos.y) * ${t.params.size[0]};
+  var subgroup_id = linear + 1;
+  subgroup_id = subgroupBroadcastFirst(subgroup_id);
+
+  // Filter out possible helper invocations.
+  let x_in_range = u32(pos.x) < (${t.params.size[0]} - 1);
+  let y_in_range = u32(pos.y) < (${t.params.size[1]} - 1);
+  let in_range = x_in_range && y_in_range;
+  let input = select(0u, inputs[linear], in_range);
+
+  let res = select(0u, 1u, subgroupAny(bool(input)));
+  return vec2u(res, subgroup_id);
+}`;
+
+  await runFragmentTest(
+    t,
+    t.params.format,
+    fsShader,
+    t.params.size[0],
+    t.params.size[1],
+    inputData,
+    (data) => {
+      return checkFragmentAny(
+        data,
+        inputData,
+        t.params.format,
+        t.params.size[0],
+        t.params.size[1]
+      );
+    }
+  );
+});
+
+// Using subgroup operations in control with fragment shaders
+// quickly leads to unportable behavior.
+g.test('fragment,split').unimplemented();

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/subgroupBallot.spec.js

@@ -0,0 +1,339 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for subgroupBallot
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow or returning early.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf } from '../../../../../../common/util/data_tables.js';
+import { iterRange } from '../../../../../../common/util/util.js';
+import { GPUTest } from '../../../../../gpu_test.js';
+
+export const g = makeTestGroup(GPUTest);
+
+// 128 is the maximum possible subgroup size.
+const kInvocations = 128;
+
+function getMask(size) {
+  return (1n << BigInt(size)) - 1n;
+}
+
+function checkBallots(
+data,
+subgroupSize,
+filter,
+expect,
+allActive)
+{
+  for (let i = 0; i < kInvocations; i++) {
+    const idx = i * 4;
+    let actual = 0n;
+    for (let j = 0; j < 4; j++) {
+      actual |= BigInt(data[idx + j]) << BigInt(32 * j);
+    }
+    let expectedResult = expect(subgroupSize);
+    const subgroupId = i % subgroupSize;
+    if (!allActive && !filter(subgroupId, subgroupSize)) {
+      expectedResult = 0n;
+    }
+    if (expectedResult !== actual) {
+      return new Error(
+        `Invocation ${i}, subgroup inv id ${i % subgroupSize}, size ${subgroupSize}
+- expected: ${expectedResult.toString(16)}
+-      got: ${actual.toString(16)}`
+      );
+    }
+  }
+
+  return undefined;
+}
+
+async function runTest(
+t,
+wgsl,
+filter,
+expect,
+allActive)
+{
+  const sizeBuffer = t.makeBufferWithContents(
+    new Uint32Array([0]),
+    GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE
+  );
+  t.trackForCleanup(sizeBuffer);
+
+  const outputNumInts = kInvocations * 4;
+  const outputBuffer = t.makeBufferWithContents(
+    new Uint32Array([...iterRange(outputNumInts, (x) => 0)]),
+    GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE
+  );
+  t.trackForCleanup(outputBuffer);
+
+  const pipeline = t.device.createComputePipeline({
+    layout: 'auto',
+    compute: {
+      module: t.device.createShaderModule({
+        code: wgsl
+      }),
+      entryPoint: 'main'
+    }
+  });
+  const bg = t.device.createBindGroup({
+    layout: pipeline.getBindGroupLayout(0),
+    entries: [
+    {
+      binding: 0,
+      resource: {
+        buffer: sizeBuffer
+      }
+    },
+    {
+      binding: 1,
+      resource: {
+        buffer: outputBuffer
+      }
+    }]
+
+  });
+
+  const encoder = t.device.createCommandEncoder();
+  const pass = encoder.beginComputePass();
+  pass.setPipeline(pipeline);
+  pass.setBindGroup(0, bg);
+  pass.dispatchWorkgroups(1, 1, 1);
+  pass.end();
+  t.queue.submit([encoder.finish()]);
+
+  const sizeReadback = await t.readGPUBufferRangeTyped(sizeBuffer, {
+    srcByteOffset: 0,
+    type: Uint32Array,
+    typedLength: 1,
+    method: 'copy'
+  });
+  const subgroupSize = sizeReadback.data[0];
+
+  const outputReadback = await t.readGPUBufferRangeTyped(outputBuffer, {
+    srcByteOffset: 0,
+    type: Uint32Array,
+    typedLength: outputNumInts,
+    method: 'copy'
+  });
+  const output = outputReadback.data;
+
+  t.expectOK(checkBallots(output, subgroupSize, filter, expect, allActive));
+}
+
+const kCases = {
+  every_even: {
+    cond: `id % 2 == 0`,
+    filter: (id, size) => {
+      return id % 2 === 0;
+    },
+    expect: (subgroupSize) => {
+      const base = BigInt('0x55555555555555555555555555555555');
+      const mask = getMask(subgroupSize);
+      return base & mask;
+    }
+  },
+  every_odd: {
+    cond: `id % 2 == 1`,
+    filter: (id, size) => {
+      return id % 2 === 1;
+    },
+    expect: (subgroupSize) => {
+      const base = BigInt('0xAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA');
+      const mask = getMask(subgroupSize);
+      return base & mask;
+    }
+  },
+  lower_half: {
+    cond: `id < subgroupSize / 2`,
+    filter: (id, size) => {
+      return id < Math.floor(size / 2);
+    },
+    expect: (size) => {
+      return getMask(Math.floor(size / 2));
+    }
+  },
+  upper_half: {
+    cond: `id >= subgroupSize / 2`,
+    filter: (id, size) => {
+      return id >= Math.floor(size / 2);
+    },
+    expect: (size) => {
+      return getMask(Math.floor(size / 2)) << BigInt(Math.floor(size / 2));
+    }
+  },
+  first_two: {
+    cond: `id == 0 || id == 1`,
+    filter: (id) => {
+      return id === 0 || id === 1;
+    },
+    expect: (size) => {
+      return getMask(2);
+    }
+  }
+};
+
+g.test('compute,split').
+desc('Tests ballot in a split subgroup').
+params((u) => u.combine('case', keysOf(kCases))).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kCases[t.params.case];
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> size : u32;
+
+@group(0) @binding(1)
+var<storage, read_write> output : array<vec4u>;
+
+@compute @workgroup_size(${kInvocations})
+fn main(@builtin(subgroup_size) subgroupSize : u32,
+        @builtin(subgroup_invocation_id) id : u32,
+        @builtin(local_invocation_index) lid : u32) {
+  if (lid == 0) {
+    size = subgroupSize;
+  }
+  if ${testcase.cond} {
+    output[lid] = subgroupBallot(true);
+  } else {
+    return;
+  }
+}`;
+
+  await runTest(t, wgsl, testcase.filter, testcase.expect, false);
+});
+
+g.test('fragment,split').unimplemented();
+
+g.test('predicate').
+desc('Tests the predicate parameter').
+params((u) => u.combine('case', keysOf(kCases))).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kCases[t.params.case];
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> size : u32;
+
+@group(0) @binding(1)
+var<storage, read_write> output : array<vec4u>;
+
+@compute @workgroup_size(${kInvocations})
+fn main(@builtin(subgroup_size) subgroupSize : u32,
+        @builtin(subgroup_invocation_id) id : u32,
+        @builtin(local_invocation_index) lid : u32) {
+  if (lid == 0) {
+    size = subgroupSize;
+  }
+  let cond = ${testcase.cond};
+  let b = subgroupBallot(cond);
+  output[lid] = b;
+}`;
+
+  await runTest(t, wgsl, testcase.filter, testcase.expect, true);
+});
+
+const kBothCases = {
+  empty: {
+    cond: `id < subgroupSize / 2`,
+    pred: `id >= subgroupSize / 2`,
+    filter: (id, size) => {
+      return id < Math.floor(size / 2);
+    },
+    expect: (size) => {
+      return 0n;
+    }
+  },
+  full: {
+    cond: `id < 128`,
+    pred: `lid < 128`,
+    filter: (id, size) => {
+      return true;
+    },
+    expect: (size) => {
+      return getMask(size);
+    }
+  },
+  one_in_four: {
+    cond: `id % 2 == 0`,
+    pred: `id % 4 == 0`,
+    filter: (id, size) => {
+      return id % 2 === 0;
+    },
+    expect: (size) => {
+      const base = BigInt('0x11111111111111111111111111111111');
+      const mask = getMask(size);
+      return base & mask;
+    }
+  },
+  middle_half: {
+    cond: `id >= subgroupSize / 4`,
+    pred: `id < 3 * (subgroupSize / 4)`,
+    filter: (id, size) => {
+      return id >= Math.floor(size / 4);
+    },
+    expect: (size) => {
+      return getMask(Math.floor(size / 2)) << BigInt(Math.floor(size / 4));
+    }
+  },
+  middle_half_every_other: {
+    cond: `(id >= subgroupSize / 4) && (id < 3 * (subgroupSize / 4))`,
+    pred: `id % 2 == 0`,
+    filter: (id, size) => {
+      return id >= Math.floor(size / 4) && id < 3 * Math.floor(size / 4);
+    },
+    expect: (size) => {
+      const base = BigInt('0x55555555555555555555555555555555');
+      const mask = getMask(Math.floor(size / 2)) << BigInt(Math.floor(size / 4));
+      return base & mask;
+    }
+  }
+};
+
+g.test('predicate_and_control_flow').
+desc('Test dynamic predicate and control flow together').
+params((u) => u.combine('case', keysOf(kBothCases))).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kBothCases[t.params.case];
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> size : u32;
+
+@group(0) @binding(1)
+var<storage, read_write> output : array<vec4u>;
+
+@compute @workgroup_size(${kInvocations})
+fn main(@builtin(subgroup_size) subgroupSize : u32,
+        @builtin(subgroup_invocation_id) id : u32,
+        @builtin(local_invocation_index) lid : u32) {
+  if (lid == 0) {
+    size = subgroupSize;
+  }
+  if ${testcase.cond} {
+    output[lid] = subgroupBallot(${testcase.pred});
+  } else {
+    return;
+  }
+}`;
+
+  await runTest(t, wgsl, testcase.filter, testcase.expect, false);
+});
+
+g.test('fragment').unimplemented();

tests/wpt/webgpu/tests/webgpu/webgpu/shader/execution/expression/call/builtin/subgroupBitwise.spec.js

@@ -0,0 +1,562 @@
+/**
+* AUTO-GENERATED - DO NOT EDIT. Source: https://github.com/gpuweb/cts
+**/export const description = `
+Execution tests for subgroupAny.
+
+Note: There is a lack of portability for non-uniform execution so these tests
+restrict themselves to uniform control flow.
+Note: There is no guaranteed mapping between subgroup_invocation_id and
+local_invocation_index. Tests should avoid assuming there is.
+`;import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
+import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js';
+import { iterRange } from '../../../../../../common/util/util.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import {
+  kConcreteSignedIntegerScalarsAndVectors,
+  kConcreteUnsignedIntegerScalarsAndVectors,
+  scalarTypeOf,
+
+  VectorType } from
+'../../../../../util/conversion.js';
+import { align } from '../../../../../util/math.js';
+import { PRNG } from '../../../../../util/prng.js';
+
+import {
+  kWGSizes,
+  kPredicateCases,
+  SubgroupTest,
+  kDataSentinel,
+  runComputeTest,
+  runFragmentTest,
+  kFramebufferSizes } from
+'./subgroup_util.js';
+
+export const g = makeTestGroup(SubgroupTest);
+
+const kNumCases = 15;
+const kOps = ['subgroupAnd', 'subgroupOr', 'subgroupXor'];
+const kTypes = objectsToRecord([
+...kConcreteSignedIntegerScalarsAndVectors,
+...kConcreteUnsignedIntegerScalarsAndVectors]
+);
+
+/**
+ * Performs the appropriate bitwise operation on v1 and v2.
+ *
+ * @param op The subgroup operation
+ * @param v1 The first value
+ * @param v2 The second value
+ */
+function bitwise(op, v1, v2) {
+  switch (op) {
+    case 'subgroupAnd':
+      return v1 & v2;
+    case 'subgroupOr':
+      return v1 | v2;
+    case 'subgroupXor':
+      return v1 ^ v2;
+  }
+}
+
+/**
+ * Returns the identity value for the subgroup operations
+ *
+ * @param op The subgroup operation
+ */
+function identity(op) {
+  switch (op) {
+    case 'subgroupAnd':
+      return ~0;
+    case 'subgroupOr':
+    case 'subgroupXor':
+      return 0;
+  }
+}
+
+/**
+ * Checks the results for data type test
+ *
+ * The shader generate a unique subgroup id for each subgroup (avoiding 0).
+ * The check calculates the expected result for all subgroups and then compares that
+ * to the actual results.
+ * @param metadata An array of integers divided as follows:
+ *                 * first half subgroup invocation id
+ *                 * second half unique subgroup id
+ * @param output An array of output values
+ * @param type The type being tested
+ * @param op The subgroup operation
+ * @param offset A constant offset added to subgroup invocation id to form the
+ *               the input to the subgroup operation
+ */
+function checkDataTypes(
+metadata,
+output,
+type,
+op,
+offset)
+{
+  const expected = new Map();
+  for (let i = 0; i < Math.floor(metadata.length / 2); i++) {
+    const group_id = metadata[i + Math.floor(metadata.length / 2)];
+    let expect = expected.get(group_id) ?? identity(op);
+    expect = bitwise(op, expect, i + offset);
+    expected.set(group_id, expect);
+  }
+
+  let numEles = 1;
+  let stride = 1;
+  if (type instanceof VectorType) {
+    numEles = type.width;
+    stride = numEles === 3 ? 4 : numEles;
+  }
+  for (let inv = 0; inv < Math.floor(output.length / stride); inv++) {
+    const group_id = metadata[inv + Math.floor(metadata.length / 2)];
+    const expect = expected.get(group_id) ?? 0;
+    for (let ele = 0; ele < numEles; ele++) {
+      const res = output[inv * stride + ele];
+      if (res !== expect) {
+        return new Error(`Invocation ${inv}, component ${ele}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('data_types').
+desc('Tests allowed data types').
+params((u) =>
+u.
+combine('type', keysOf(kTypes)).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('op', kOps)
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const type = kTypes[t.params.type];
+  let numEles = 1;
+  if (type instanceof VectorType) {
+    numEles = type.width === 3 ? 4 : type.width;
+  }
+
+  const scalarTy = scalarTypeOf(type);
+
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<${type.toString()}>;
+
+struct Metadata {
+  id : array<u32, ${wgThreads}>,
+  group_id : array<u32, ${wgThreads}>
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+) {
+
+  // Record subgroup invocation id for this invocation.
+  metadata.id[lid] = id;
+
+  // Record a unique id for this subgroup (avoid 0).
+  let group_id = subgroupBroadcastFirst(lid + 1);
+  metadata.group_id[lid] = group_id;
+
+  outputs[lid] = ${t.params.op}(${type.toString()}(${scalarTy.toString()}(lid + inputs[0])));
+}`;
+
+  const magicOffset = 0x7fff000f;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    numEles,
+    new Uint32Array([magicOffset]),
+    (metadata, output) => {
+      return checkDataTypes(metadata, output, type, t.params.op, magicOffset);
+    }
+  );
+});
+
+/**
+ * Generates randomized input data
+ *
+ * Case 0: All 0s
+ * Case 1: All 0xffffs
+ * Case 2-9: All identity values except an inverted value randomly every 32 values.
+ *           All values capped to 0xffff
+ * Case 10+: Random values in the range [0, 2 ** 30]
+ * @param seed The PRNG seed
+ * @param num The number of values to generate
+ * @param identity The identity value for the operation
+ */
+function generateInputData(seed, num, identity) {
+  const prng = new PRNG(seed);
+
+  const bound = Math.min(num, 32);
+  const index = prng.uniformInt(bound);
+
+  return new Uint32Array([
+  ...iterRange(num, (x) => {
+    if (seed === 0) {
+      return 0;
+    } else if (seed === 1) {
+      return 0xffff;
+    } else if (seed < 10) {
+      const bounded = x % bound;
+      let val = bounded === index ? ~identity : identity;
+      val &= 0xffff;
+      return val;
+    }
+    return prng.uniformInt(1 << 30);
+  })]
+  );
+}
+
+/**
+ * Checks the result of compute tests
+ *
+ * Calculates the expected results for each subgroup and compares against
+ * the actual output.
+ * @param metadata An array divided as follows:
+ *                 * first half: subgroup invocation id in lower 16 bits
+ *                               subgroup size in upper 16 bits
+ *                 * second half: unique subgroup id
+ * @param output The outputs
+ * @param input The input data
+ * @param op The subgroup operation
+ * @param filter A predicate used to filter invocations.
+ */
+function checkBitwiseCompute(
+metadata,
+output,
+input,
+op,
+filter)
+{
+  const expected = new Map();
+  for (let i = 0; i < output.length; i++) {
+    const group_id = metadata[i + output.length];
+    const combo = metadata[i];
+    const id = combo & 0xffff;
+    const size = combo >> 16 & 0xffff;
+    if (filter(id, size)) {
+      let expect = expected.get(group_id) ?? identity(op);
+      expect = bitwise(op, expect, input[i]);
+      expected.set(group_id, expect);
+    }
+  }
+
+  for (let i = 0; i < output.length; i++) {
+    const group_id = metadata[i + output.length];
+    const combo = metadata[i];
+    const id = combo & 0xffff;
+    const size = combo >> 16 & 0xffff;
+    const res = output[i];
+    if (filter(id, size)) {
+      const expect = expected.get(group_id) ?? 0;
+      if (res !== expect) {
+        return new Error(`Invocation ${i}: incorrect result
+- expected: ${expect}
+-      got: ${res}`);
+      }
+    } else {
+      if (res !== kDataSentinel) {
+        return new Error(`Invocation ${i}: unexpected write`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('compute,all_active').
+desc('Test bitwise operations with randomized inputs').
+params((u) =>
+u.
+combine('case', [...iterRange(kNumCases, (x) => x)]).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('op', kOps)
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  id_and_size : array<u32, ${wgThreads}>,
+  group_id : array<u32, ${wgThreads}>
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) sg_size : u32,
+) {
+
+  // Record both subgroup invocation id and subgroup size in the same u32.
+  // Subgroups sizes are in the range [4, 128] so both values fit.
+  metadata.id_and_size[lid] = id | (sg_size << 16);
+
+  // Record a unique id for this subgroup (avoid 0).
+  let group_id = subgroupBroadcastFirst(lid + 1);
+  metadata.group_id[lid] = group_id;
+
+  outputs[lid] = ${t.params.op}(inputs[lid]);
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads, identity(t.params.op));
+  const uintsPerOutput = 1;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkBitwiseCompute(
+        metadata,
+        output,
+        inputData,
+        t.params.op,
+        (id, size) => {
+          return true;
+        }
+      );
+    }
+  );
+});
+
+g.test('compute,split').
+desc('Test that only active invocations participate').
+params((u) =>
+u.
+combine('predicate', keysOf(kPredicateCases)).
+beginSubcases().
+combine('wgSize', kWGSizes).
+combine('op', kOps).
+combine('case', [...iterRange(kNumCases, (x) => x)])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const testcase = kPredicateCases[t.params.predicate];
+  const wgThreads = t.params.wgSize[0] * t.params.wgSize[1] * t.params.wgSize[2];
+
+  const wgsl = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage> inputs : array<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> outputs : array<u32>;
+
+struct Metadata {
+  id_and_size : array<u32, ${wgThreads}>,
+  group_id : array<u32, ${wgThreads}>
+}
+
+@group(0) @binding(2)
+var<storage, read_write> metadata : Metadata;
+
+@compute @workgroup_size(${t.params.wgSize[0]}, ${t.params.wgSize[1]}, ${t.params.wgSize[2]})
+fn main(
+  @builtin(local_invocation_index) lid : u32,
+  @builtin(subgroup_invocation_id) id : u32,
+  @builtin(subgroup_size) subgroupSize : u32,
+) {
+
+  // Record both subgroup invocation id and subgroup size in the same u32.
+  // Subgroups sizes are in the range [4, 128] so both values fit.
+  metadata.id_and_size[lid] = id | (subgroupSize << 16);
+
+  // Record a unique id for this subgroup (avoid 0).
+  let group_id = subgroupBroadcastFirst(lid + 1);
+  metadata.group_id[lid] = group_id;
+
+  if ${testcase.cond} {
+    outputs[lid] = ${t.params.op}(inputs[lid]);
+  } else {
+    return;
+  }
+}`;
+
+  const inputData = generateInputData(t.params.case, wgThreads, identity(t.params.op));
+  const uintsPerOutput = 1;
+  await runComputeTest(
+    t,
+    wgsl,
+    [t.params.wgSize[0], t.params.wgSize[1], t.params.wgSize[2]],
+    uintsPerOutput,
+    inputData,
+    (metadata, output) => {
+      return checkBitwiseCompute(metadata, output, inputData, t.params.op, testcase.filter);
+    }
+  );
+});
+
+/**
+ * Checks bitwise ops results from a fragment shader.
+ *
+ * Avoids the last row and column to skip potential helper invocations.
+ * @param data Framebuffer output
+ *             * component 0 is result
+ *             * component 1 is generated subgroup id
+ * @param input An array of input data
+ * @param op The subgroup operation
+ * @param format The framebuffer format
+ * @param width Framebuffer width
+ * @param height Framebuffer height
+ */
+function checkBitwiseFragment(
+data,
+input,
+op,
+format,
+width,
+height)
+{
+  const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format];
+  const blocksPerRow = width / blockWidth;
+  // 256 minimum comes from image copy requirements.
+  const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256);
+  const uintsPerRow = bytesPerRow / 4;
+  const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4;
+
+  // Iteration skips last row and column to avoid helper invocations because it is not
+  // guaranteed whether or not they participate in the subgroup operation.
+  const expected = new Map();
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        return new Error(`Internal error: helper invocation at (${col}, ${row})`);
+      }
+
+      let v = expected.get(subgroup_id) ?? identity(op);
+      v = bitwise(op, v, input[row * width + col]);
+      expected.set(subgroup_id, v);
+    }
+  }
+
+  for (let row = 0; row < height - 1; row++) {
+    for (let col = 0; col < width - 1; col++) {
+      const offset = uintsPerRow * row + col * uintsPerTexel;
+      const res = data[offset];
+      const subgroup_id = data[offset + 1];
+
+      if (subgroup_id === 0) {
+        // Inactive in the fragment.
+        continue;
+      }
+
+      const expected_v = expected.get(subgroup_id) ?? 0;
+      if (expected_v !== res) {
+        return new Error(`Row ${row}, col ${col}: incorrect results:
+- expected: ${expected_v}
+-      got: ${res}`);
+      }
+    }
+  }
+
+  return undefined;
+}
+
+g.test('fragment,all_active').
+desc('Tests bitwise operations in fragment shaders').
+params((u) =>
+u.
+combine('size', kFramebufferSizes).
+beginSubcases().
+combine('case', [...iterRange(kNumCases, (x) => x)]).
+combine('op', kOps).
+combineWithParams([{ format: 'rg32uint' }])
+).
+beforeAllSubcases((t) => {
+  t.selectDeviceOrSkipTestCase('subgroups');
+}).
+fn(async (t) => {
+  const numInputs = t.params.size[0] * t.params.size[1];
+  const inputData = generateInputData(t.params.case, numInputs, identity(t.params.op));
+
+  const ident = identity(t.params.op) === 0 ? '0' : '0xffffffff';
+  const fsShader = `
+enable subgroups;
+
+@group(0) @binding(0)
+var<storage, read_write> inputs : array<u32>;
+
+@fragment
+fn main(
+  @builtin(position) pos : vec4f,
+) -> @location(0) vec2u {
+  // Generate a subgroup id based on linearized position, avoid 0.
+  let linear = u32(pos.x) + u32(pos.y) * ${t.params.size[0]};
+  let subgroup_id = subgroupBroadcastFirst(linear + 1);
+
+  // Filter out possible helper invocations.
+  let x_in_range = u32(pos.x) < (${t.params.size[0]} - 1);
+  let y_in_range = u32(pos.y) < (${t.params.size[1]} - 1);
+  let in_range = x_in_range && y_in_range;
+  let input = select(${ident}, inputs[linear], in_range);
+
+  let res = ${t.params.op}(input);
+  return vec2u(res, subgroup_id);
+}`;
+
+  await runFragmentTest(
+    t,
+    t.params.format,
+    fsShader,
+    t.params.size[0],
+    t.params.size[1],
+    inputData,
+    (data) => {
+      return checkBitwiseFragment(
+        data,
+        inputData,
+        t.params.op,
+        t.params.format,
+        t.params.size[0],
+        t.params.size[1]
+      );
+    }
+  );
+});
+
+g.test('fragment,split').unimplemented();