WGSL textureNumLayers execution test (gpuweb#3853)

src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts

@@ -5,9 +5,35 @@ Returns the number of layers (elements) of an array texture.
 `;
 
 import { makeTestGroup } from '../../../../../../common/framework/test_group.js';
-import { GPUTest } from '../../../../../gpu_test.js';
+import { kTextureFormatInfo } from '../../../../../format_info.js';
+import { TexelFormats } from '../../../../types.js';
 
-export const g = makeTestGroup(GPUTest);
+import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js';
+
+const kNumLayers = 36;
+
+function getLayerSettingsAndExpected({
+  view_type,
+  isCubeArray,
+}: {
+  view_type: 'full' | 'partial';
+  isCubeArray?: boolean;
+}) {
+  const divisor = isCubeArray ? 6 : 1;
+  return view_type === 'partial'
+    ? {
+        baseArrayLayer: 11,
+        arrayLayerCount: 6,
+        expected: [6 / divisor],
+      }
+    : {
+        baseArrayLayer: 0,
+        arrayLayerCount: kNumLayers,
+        expected: [kNumLayers / divisor],
+      };
+}
+
+export const g = makeTestGroup(WGSLTextureQueryTest);
 
 g.test('sampled')
   .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers')
@@ -26,9 +52,49 @@ Parameters
     u
       .combine('texture_type', ['texture_2d_array', 'texture_cube_array'] as const)
       .beginSubcases()
-      .combine('sampled_type', ['f32-only', 'i32', 'u32'] as const)
+      .combine('sampled_type', ['f32', 'i32', 'u32'] as const)
+      .combine('view_type', ['full', 'partial'] as const)
   )
-  .unimplemented();
+  .beforeAllSubcases(t => {
+    t.skipIf(
+      t.isCompatibility && t.params.view === 'partial',
+      'compatibility mode does not support partial layer views'
+    );
+    t.skipIf(
+      t.isCompatibility && t.params.texture_type === 'texture_cube_array',
+      'compatibility mode does not support cube arrays'
+    );
+  })
+  .fn(t => {
+    const { texture_type, sampled_type, view_type } = t.params;
+    const { format } = kSampleTypeInfo[sampled_type];
+
+    const texture = t.createTextureTracked({
+      format,
+      usage: GPUTextureUsage.TEXTURE_BINDING,
+      size: [1, 1, kNumLayers],
+    });
+
+    const code = `
+@group(0) @binding(0) var t: ${texture_type}<${sampled_type}>;
+@group(0) @binding(1) var<storage, read_write> result: u32;
+@compute @workgroup_size(1) fn cs() {
+  result = textureNumLayers(t);
+}
+    `;
+
+    const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({
+      view_type,
+      isCubeArray: texture_type === 'texture_cube_array',
+    });
+    const view = texture.createView({
+      dimension: texture_type === 'texture_2d_array' ? '2d-array' : 'cube-array',
+      baseArrayLayer,
+      arrayLayerCount,
+    });
+
+    t.executeAndExpectResult(code, view, expected);
+  });
 
 g.test('arrayed')
   .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers')
@@ -42,9 +108,50 @@ Parameters
 `
   )
   .params(u =>
-    u.combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const)
+    u
+      .combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const)
+      .beginSubcases()
+      .combine('view_type', ['full', 'partial'] as const)
   )
-  .unimplemented();
+  .beforeAllSubcases(t => {
+    t.skipIf(
+      t.isCompatibility && t.params.view === 'partial',
+      'compatibility mode does not support partial layer views'
+    );
+    t.skipIf(
+      t.isCompatibility && t.params.texture_type === 'texture_depth_cube_array',
+      'compatibility mode does not support cube arrays'
+    );
+  })
+  .fn(t => {
+    const { texture_type, view_type } = t.params;
+
+    const texture = t.createTextureTracked({
+      format: 'depth32float',
+      usage: GPUTextureUsage.TEXTURE_BINDING,
+      size: [1, 1, kNumLayers],
+    });
+
+    const code = `
+@group(0) @binding(0) var t: ${texture_type};
+@group(0) @binding(1) var<storage, read_write> result: u32;
+@compute @workgroup_size(1) fn cs() {
+  result = textureNumLayers(t);
+}
+    `;
+
+    const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({
+      view_type,
+      isCubeArray: texture_type === 'texture_depth_cube_array',
+    });
+    const view = texture.createView({
+      dimension: texture_type === 'texture_depth_2d_array' ? '2d-array' : 'cube-array',
+      baseArrayLayer,
+      arrayLayerCount,
+    });
+
+    t.executeAndExpectResult(code, view, expected);
+  });
 
 g.test('storage')
   .specURL('https://www.w3.org/TR/WGSL/#texturenumlayers')
@@ -76,25 +183,40 @@ Parameters
   )
   .params(u =>
     u
+      .combineWithParams(TexelFormats)
       .beginSubcases()
-      .combine('texel_format', [
-        'rgba8unorm',
-        'rgba8snorm',
-        'rgba8uint',
-        'rgba8sint',
-        'rgba16uint',
-        'rgba16sint',
-        'rgba16float',
-        'r32uint',
-        'r32sint',
-        'r32float',
-        'rg32uint',
-        'rg32sint',
-        'rg32float',
-        'rgba32uint',
-        'rgba32sint',
-        'rgba32float',
-      ] as const)
       .combine('access_mode', ['read', 'write', 'read_write'] as const)
+      .filter(
+        t => t.access_mode !== 'read_write' || kTextureFormatInfo[t.format].color?.readWriteStorage
+      )
+      .combine('view_type', ['full', 'partial'] as const)
   )
-  .unimplemented();
+  .beforeAllSubcases(t => t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format))
+  .fn(t => {
+    const { format, access_mode, view_type } = t.params;
+
+    const texture = t.createTextureTracked({
+      format,
+      usage: GPUTextureUsage.STORAGE_BINDING,
+      size: [1, 1, kNumLayers],
+    });
+
+    const code = `
+@group(0) @binding(0) var t: texture_storage_2d_array<${format}, ${access_mode}>;
+@group(0) @binding(1) var<storage, read_write> result: u32;
+@compute @workgroup_size(1) fn cs() {
+  result = textureNumLayers(t);
+}
+    `;
+
+    const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({
+      view_type,
+    });
+    const view = texture.createView({
+      dimension: '2d-array',
+      baseArrayLayer,
+      arrayLayerCount,
+    });
+
+    t.executeAndExpectResult(code, view, expected);
+  });

src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts

@@ -31,6 +31,61 @@ import { TexelView } from '../../../../../util/texture/texel_view.js';
 import { createTextureFromTexelViews } from '../../../../../util/texture.js';
 import { reifyExtent3D } from '../../../../../util/unions.js';
 
+export type SampledType = 'f32' | 'i32' | 'u32';
+
+export const kSampleTypeInfo = {
+  f32: {
+    format: 'rgba8unorm',
+  },
+  i32: {
+    format: 'rgba32sint',
+  },
+  u32: {
+    format: 'rgba32uint',
+  },
+} as const;
+
+/**
+ * Used for textureDimension, textureNumLevels, textureNumLayers
+ */
+export class WGSLTextureQueryTest extends GPUTest {
+  executeAndExpectResult(code: string, view: GPUTextureView, expected: number[]) {
+    const { device } = this;
+    const module = device.createShaderModule({ code });
+    const pipeline = device.createComputePipeline({
+      layout: 'auto',
+      compute: {
+        module,
+      },
+    });
+
+    const resultBuffer = this.createBufferTracked({
+      size: 16,
+      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
+    });
+
+    const bindGroup = device.createBindGroup({
+      layout: pipeline.getBindGroupLayout(0),
+      entries: [
+        { binding: 0, resource: view },
+        { binding: 1, resource: { buffer: resultBuffer } },
+      ],
+    });
+
+    const encoder = device.createCommandEncoder();
+    const pass = encoder.beginComputePass();
+    pass.setPipeline(pipeline);
+    pass.setBindGroup(0, bindGroup);
+    pass.dispatchWorkgroups(1);
+    pass.end();
+    device.queue.submit([encoder.finish()]);
+
+    const e = new Uint32Array(4);
+    e.set(expected);
+    this.expectGPUBufferValuesEqual(resultBuffer, e);
+  }
+}
+
 function getLimitValue(v: number) {
   switch (v) {
     case Number.POSITIVE_INFINITY: