shader/validation: Validate binding points on variables (gpuweb#3348)

Test that:
- resource variables have both @group and @binding
- non-resource variables do not have @group or @binding
- binding points cannot collide in the same entry point
- binding points *can* collide in different entry points

Remove some tests from other suites that are now better covered here.

Issue gpuweb#1570

src/webgpu/listing_meta.json

@@ -1733,6 +1733,11 @@
   "webgpu:shader,validation,decl,ptr_spelling:ptr_bad_store_type:*": { "subcaseMS": 0.967 },
   "webgpu:shader,validation,decl,ptr_spelling:ptr_handle_space_invalid:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,ptr_spelling:ptr_not_instantiable:*": { "subcaseMS": 1.310 },
+  "webgpu:shader,validation,decl,var:binding_point_on_resources:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_point_on_non_resources:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_point_on_function_var:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_collisions:*": { "subcaseMS": 1.000 },
+  "webgpu:shader,validation,decl,var:binding_collision_unused_helper:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:module_scope_types:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:function_scope_types:*": { "subcaseMS": 1.000 },
   "webgpu:shader,validation,decl,var:function_addrspace_at_module_scope:*": { "subcaseMS": 1.000 },
@@ -1989,7 +1994,6 @@
   "webgpu:shader,validation,parse,var_and_let:var_access_mode_bad_template_delim:*": { "subcaseMS": 1.088 },
   "webgpu:shader,validation,shader_io,binding:binding:*": { "subcaseMS": 1.240 },
   "webgpu:shader,validation,shader_io,binding:binding_f16:*": { "subcaseMS": 0.500 },
-  "webgpu:shader,validation,shader_io,binding:binding_without_group:*": { "subcaseMS": 0.901 },
   "webgpu:shader,validation,shader_io,builtins:duplicates:*": { "subcaseMS": 1.913 },
   "webgpu:shader,validation,shader_io,builtins:missing_vertex_position:*": { "subcaseMS": 0.975 },
   "webgpu:shader,validation,shader_io,builtins:nesting:*": { "subcaseMS": 2.700 },
@@ -2003,7 +2007,6 @@
   "webgpu:shader,validation,shader_io,entry_point:no_entry_point_provided:*": { "subcaseMS": 0.801 },
   "webgpu:shader,validation,shader_io,group:group:*": { "subcaseMS": 1.355 },
   "webgpu:shader,validation,shader_io,group:group_f16:*": { "subcaseMS": 0.400 },
-  "webgpu:shader,validation,shader_io,group:group_without_binding:*": { "subcaseMS": 1.100 },
   "webgpu:shader,validation,shader_io,group_and_binding:binding_attributes:*": { "subcaseMS": 1.280 },
   "webgpu:shader,validation,shader_io,group_and_binding:different_entry_points:*": { "subcaseMS": 1.833 },
   "webgpu:shader,validation,shader_io,group_and_binding:function_scope:*": { "subcaseMS": 1.000 },

src/webgpu/shader/validation/decl/var.spec.ts

@@ -338,3 +338,102 @@ g.test('function_addrspace_at_module_scope')
       `var<${t.params.addrspace}> foo : i32;`
     );
   });
+
+// A list of resource variable declarations.
+const kResourceDecls = {
+  uniform: 'var<uniform> buffer : vec4f;',
+  storage: 'var<storage> buffer : vec4f;',
+  texture: 'var t : texture_2d<f32>;',
+  sampler: 'var s : sampler;',
+};
+
+g.test('binding_point_on_resources')
+  .desc('Test that resource variables must have both @group and @binding attributes.')
+  .params(u =>
+    u
+      .combine('decl', keysOf(kResourceDecls))
+      .combine('group', ['', '@group(0)'])
+      .combine('binding', ['', '@binding(0)'])
+  )
+  .fn(t => {
+    const shouldPass = t.params.group !== '' && t.params.binding !== '';
+    const wgsl = `${t.params.group} ${t.params.binding} ${kResourceDecls[t.params.decl]}`;
+    t.expectCompileResult(shouldPass, wgsl);
+  });
+
+g.test('binding_point_on_non_resources')
+  .desc('Test that non-resource variables cannot have either @group or @binding attributes.')
+  .params(u =>
+    u
+      .combine('addrspace', ['private', 'workgroup'])
+      .combine('group', ['', '@group(0)'])
+      .combine('binding', ['', '@binding(0)'])
+  )
+  .fn(t => {
+    const shouldPass = t.params.group === '' && t.params.binding === '';
+    const wgsl = `${t.params.group} ${t.params.binding} var<${t.params.addrspace}> foo : i32;`;
+    t.expectCompileResult(shouldPass, wgsl);
+  });
+
+g.test('binding_point_on_function_var')
+  .desc('Test that function variables cannot have either @group or @binding attributes.')
+  .params(u => u.combine('group', ['', '@group(0)']).combine('binding', ['', '@binding(0)']))
+  .fn(t => {
+    const shouldPass = t.params.group === '' && t.params.binding === '';
+    const wgsl = `
+    fn foo() {
+      ${t.params.group} ${t.params.binding} var bar : i32;
+    }`;
+    t.expectCompileResult(shouldPass, wgsl);
+  });
+
+g.test('binding_collisions')
+  .desc('Test that binding points can collide iff they are not used by the same entry point.')
+  .params(u =>
+    u
+      .combine('a_group', [0, 1])
+      .combine('b_group', [0, 1])
+      .combine('a_binding', [0, 1])
+      .combine('b_binding', [0, 1])
+      .combine('b_use', ['same', 'different'])
+  )
+  .fn(t => {
+    const wgsl = `
+    @group(${t.params.a_group}) @binding(${t.params.a_binding}) var<uniform> a : vec4f;
+    @group(${t.params.b_group}) @binding(${t.params.b_binding}) var<uniform> b : vec4f;
+
+    @fragment
+    fn main1() {
+      _ = a;
+      ${
+        t.params.b_use === 'same'
+          ? ''
+          : `
+      }
+
+    @fragment
+    fn main2() {`
+      }
+      _ = b;
+    }`;
+
+    const collision =
+      t.params.a_group === t.params.b_group && t.params.a_binding === t.params.b_binding;
+    const shouldFail = collision && t.params.b_use === 'same';
+    t.expectCompileResult(!shouldFail, wgsl);
+  });
+
+g.test('binding_collision_unused_helper')
+  .desc('Test that binding points can collide in an unused helper function.')
+  .fn(t => {
+    const wgsl = `
+    @group(0) @binding(0) var<uniform> a : vec4f;
+    @group(0) @binding(0) var<uniform> b : vec4f;
+
+    fn foo() {
+      _ = a;
+      _ = b;
+    }`;
+
+    t.expectCompileResult(true, wgsl);
+  });

src/webgpu/shader/validation/shader_io/binding.spec.ts

@@ -124,17 +124,3 @@ var<storage> a: i32;
 }`;
     t.expectCompileResult(false, code);
   });
-
-g.test('binding_without_group')
-  .desc(`Test validation of binding without group`)
-  .fn(t => {
-    const code = `
-@binding(1)
-var<storage> a: i32;
-
-@workgroup_size(1, 1, 1)
-@compute fn main() {
-  _ = a;
-}`;
-    t.expectCompileResult(false, code);
-  });

src/webgpu/shader/validation/shader_io/group.spec.ts

@@ -124,17 +124,3 @@ var<storage> a: i32;
 }`;
     t.expectCompileResult(false, code);
   });
-
-g.test('group_without_binding')
-  .desc(`Test validation of group without binding`)
-  .fn(t => {
-    const code = `
-@group(1)
-var<storage> a: i32;
-
-@workgroup_size(1, 1, 1)
-@compute fn main() {
-  _ = a;
-}`;
-    t.expectCompileResult(false, code);
-  });