Add shader execution tests on the feature subgroup-size-control

src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts

@@ -1168,3 +1168,132 @@ fn main(@builtin(local_invocation_id) local_id : vec3u,
 
     t.expectOK(checkNumSubgroupsConsistency(countData, outputData, wgThreads, t.params.numWGs));
   });
+
+/**
+ * Returns all valid subgroup sizes for the given adapter info, i.e. all power-of-two values
+ * between subgroupMinSize and subgroupMaxSize inclusive.
+ */
+async function getValidSubgroupSizes(device: GPUDevice): Promise<number[]> {
+  interface SubgroupProperties extends GPUAdapterInfo {
+    subgroupMinSize: number;
+    subgroupMaxSize: number;
+  }
+  const { subgroupMinSize, subgroupMaxSize } = device.adapterInfo as SubgroupProperties;
+
+  const sizes: number[] = [];
+  for (let subgroupSize = subgroupMinSize; subgroupSize <= subgroupMaxSize; subgroupSize *= 2) {
+    const wgsl = `
+enable subgroups;
+enable subgroup_size_control;
+
+@compute @workgroup_size(${subgroupSize}, 1, 1) @subgroup_size(${subgroupSize})
+fn main(@builtin(local_invocation_index) lid : u32) {
+}`;
+    device.pushErrorScope('validation');
+    const module = device.createShaderModule({ code: wgsl });
+    device.createComputePipeline({
+      layout: 'auto',
+      compute: { module, entryPoint: 'main' },
+    });
+    const error = await device.popErrorScope();
+    if (error) {
+      continue;
+    }
+    sizes.push(subgroupSize);
+  }
+  return sizes;
+}
+
+g.test('subgroup_size_attribute')
+  .desc(
+    `Tests that at least one power-of-two value in [subgroupMinSize, subgroupMaxSize] can be used as
+    the @subgroup_size attribute in a simple compute pipeline.the value of the subgroup_size builtin
+    must equal the value of the @subgroup_size attribute.`
+  )
+  .params(u =>
+    u.combine('numWorkGroups', [1, 2] as const).combine('numSubgroups', [1, 2, 4] as const)
+  )
+  .fn(async t => {
+    t.skipIfDeviceDoesNotHaveFeature('subgroup-size-control' as GPUFeatureName);
+
+    const { numWorkGroups, numSubgroups } = t.params;
+
+    const subgroupSizes = await getValidSubgroupSizes(t.device);
+    t.expect(
+      subgroupSizes.length > 0,
+      `No valid @subgroup_size value found in [subgroupMinSize, subgroupMaxSize]`
+    );
+
+    for (const subgroupSize of subgroupSizes) {
+      const wgx = subgroupSize * numSubgroups;
+
+      const wgsl = `
+enable subgroups;
+enable subgroup_size_control;
+
+@group(0) @binding(0)
+var<storage, read_write> output : array<u32>;
+
+@compute @workgroup_size(${wgx}, 1, 1) @subgroup_size(${subgroupSize})
+fn main(@builtin(subgroup_size) builtin_size : u32,
+        @builtin(local_invocation_index) lid : u32,
+        @builtin(workgroup_id) wgid : vec3u) {
+  let gid = lid + wgid.x * ${wgx}u;
+  // Store 1 if builtin subgroup_size matches the @subgroup_size attribute, 0 otherwise.
+  output[gid] = select(0u, 1u, builtin_size == ${subgroupSize}u);
+}`;
+
+      const numInvocations = wgx * numWorkGroups;
+      const outputBuffer = t.makeBufferWithContents(
+        new Uint32Array([...iterRange(numInvocations, x => 0)]),
+        GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST
+      );
+      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: outputBuffer,
+            },
+          },
+        ],
+      });
+
+      const encoder = t.device.createCommandEncoder();
+      const pass = encoder.beginComputePass();
+      pass.setPipeline(pipeline);
+      pass.setBindGroup(0, bg);
+      pass.dispatchWorkgroups(numWorkGroups, 1, 1);
+      pass.end();
+      t.queue.submit([encoder.finish()]);
+
+      const outputReadback = await t.readGPUBufferRangeTyped(outputBuffer, {
+        srcByteOffset: 0,
+        type: Uint32Array,
+        typedLength: numInvocations,
+        method: 'copy',
+      });
+      const outputData: Uint32Array = outputReadback.data;
+
+      for (let i = 0; i < numInvocations; i++) {
+        if (outputData[i] !== 1) {
+          t.fail(
+            `@subgroup_size(${subgroupSize}): invocation ${i} has builtin subgroup_size != ${subgroupSize}`
+          );
+          break;
+        }
+      }
+    }
+  });