diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts index b8453011614a..ca7ae3d0655c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts +++ b/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 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 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 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); + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts index 815761a709d5..b441c110bd67 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/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: