diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts index b807ca97708b..b80c8334a55b 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureDimensions.spec.ts @@ -15,10 +15,12 @@ import { sampleTypeForFormatAndAspect, textureDimensionAndFormatCompatible, } from '../../../../../format_info.js'; -import { GPUTest } from '../../../../../gpu_test.js'; import { align } from '../../../../../util/math.js'; +import { kShaderStages, ShaderStage } from '../../../../validation/decl/util.js'; -export const g = makeTestGroup(GPUTest); +import { WGSLTextureQueryTest } from './texture_utils.js'; + +export const g = makeTestGroup(WGSLTextureQueryTest); /// The maximum number of texture mipmap levels to test. /// Keep this small to reduce memory and test permutations. @@ -218,8 +220,10 @@ function testValues(params: { * `values.expected`. */ function run( - t: GPUTest, - view: GPUTextureView, + t: WGSLTextureQueryTest, + stage: ShaderStage, + texture: GPUTexture | GPUExternalTexture, + viewDescriptor: GPUTextureViewDescriptor | undefined, textureType: string, levelArg: number | undefined, values: TestValues @@ -227,44 +231,16 @@ function run( const outputType = values.expected.length > 1 ? `vec${values.expected.length}u` : 'u32'; const wgsl = ` @group(0) @binding(0) var texture : ${textureType}; -@group(0) @binding(1) var output : ${outputType}; -@compute @workgroup_size(1) -fn main() { -output = ${ +fn getValue() -> ${outputType} { + return ${ levelArg !== undefined ? `textureDimensions(texture, ${levelArg})` : 'textureDimensions(texture)' }; } `; - const module = t.device.createShaderModule({ - code: wgsl, - }); - const pipeline = t.device.createComputePipeline({ - compute: { module }, - layout: 'auto', - }); - const outputBuffer = t.createBufferTracked({ - size: 32, - usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, - }); - const bindgroup = t.device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { binding: 0, resource: view }, - { binding: 1, 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.device.queue.submit([encoder.finish()]); - - t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array(values.expected)); + t.executeAndExpectResult(stage, wgsl, texture, viewDescriptor, values.expected); } /** @returns true if the GPUTextureViewDimension is valid for a storage texture */ @@ -314,6 +290,7 @@ Parameters: .expand('aspect', u => aspectsForFormat(u.format)) .expand('samples', u => samplesForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) .expand('dimensions', viewDimensions) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -339,11 +316,11 @@ Parameters: sampleCount: t.params.samples, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslSampledTextureType(): string { const base = t.params.samples !== 1 ? 'texture_multisampled' : 'texture'; @@ -362,7 +339,15 @@ Parameters: } } - run(t, textureView, wgslSampledTextureType(), t.params.textureDimensionsLevel, values); + run( + t, + t.params.stage, + texture, + viewDescriptor, + wgslSampledTextureType(), + t.params.textureDimensionsLevel, + values + ); }); g.test('depth') @@ -394,6 +379,7 @@ Parameters: .unless(u => u.aspect === 'stencil-only') .expand('samples', u => samplesForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) .expand('dimensions', viewDimensions) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -419,11 +405,11 @@ Parameters: sampleCount: t.params.samples, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslDepthTextureType(): string { const base = t.params.samples !== 1 ? 'texture_depth_multisampled' : 'texture_depth'; @@ -431,7 +417,15 @@ Parameters: return `${base}_${dimensions}`; } - run(t, textureView, wgslDepthTextureType(), t.params.textureDimensionsLevel, values); + run( + t, + t.params.stage, + texture, + viewDescriptor, + wgslDepthTextureType(), + t.params.textureDimensionsLevel, + values + ); }); g.test('storage') @@ -471,6 +465,15 @@ Parameters: .filter(p => kTextureFormatInfo[p.format].color?.storage === true) .expand('aspect', u => aspectsForFormat(u.format)) .beginSubcases() + .combine('stage', kShaderStages) + .combine('access', ['read', 'write', 'read_write'] as const) + // vertex stage can not use writable storage. + .unless(t => t.stage === 'vertex' && t.access !== 'read') + // Only some formats support write + .unless( + t => + kTextureFormatInfo[t.format].color.readWriteStorage === false && t.access === 'read_write' + ) .expand('dimensions', u => viewDimensions(u).filter(dimensionsValidForStorage)) .expand('textureMipCount', textureMipCount) .expand('baseMipLevel', baseMipLevel) @@ -490,19 +493,19 @@ Parameters: format: t.params.format, mipLevelCount: t.params.textureMipCount, }); - const textureView = texture.createView({ + const viewDescriptor: GPUTextureViewDescriptor = { dimension: t.params.dimensions, aspect: t.params.aspect, mipLevelCount: 1, baseMipLevel: t.params.baseMipLevel, - }); + }; function wgslStorageTextureType(): string { const dimensions = t.params.dimensions.replace('-', '_'); - return `texture_storage_${dimensions}<${t.params.format}, write>`; + return `texture_storage_${dimensions}<${t.params.format}, ${t.params.access}>`; } - run(t, textureView, wgslStorageTextureType(), undefined, values); + run(t, t.params.stage, texture, viewDescriptor, wgslStorageTextureType(), undefined, values); }); g.test('external') @@ -515,4 +518,23 @@ Parameters: * t: the external texture ` ) - .unimplemented(); + .params(u => + u + .beginSubcases() + .combine('stage', kShaderStages) + .combine('width', [8, 16, 24] as const) + .combine('height', [8, 16, 24] as const) + ) + .fn(t => { + const { stage, width, height } = t.params; + const canvas = new OffscreenCanvas(width, height); + // We have to make a context for VideoFrame to accept the canvas. + canvas.getContext('2d'); + const videoFrame = new VideoFrame(canvas, { timestamp: 0 }); + const texture = t.device.importExternalTexture({ source: videoFrame }); + + run(t, stage, texture, undefined, 'texture_external', undefined, { + size: [width, height], + expected: [width, height], + }); + }); 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 dd5d5923172f..0cb94b798a5d 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts @@ -88,13 +88,13 @@ fn getValue() -> u32 { view_type, isCubeArray: texture_type === 'texture_cube_array', }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: texture_type === 'texture_2d_array' ? '2d-array' : 'cube-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('arrayed') @@ -146,13 +146,13 @@ fn getValue() -> u32 { view_type, isCubeArray: texture_type === 'texture_depth_cube_array', }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: texture_type === 'texture_depth_2d_array' ? '2d-array' : 'cube-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('storage') @@ -223,11 +223,11 @@ fn getValue() -> u32 { const { baseArrayLayer, arrayLayerCount, expected } = getLayerSettingsAndExpected({ view_type, }); - const view = texture.createView({ + const viewDescription: GPUTextureViewDescriptor = { dimension: '2d-array', baseArrayLayer, arrayLayerCount, - }); + }; - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts index 61fb1f745646..cc509233399c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts @@ -112,13 +112,13 @@ fn getValue() -> u32 { view_type, mipCount ); - const view = texture.createView({ + const viewDescription = { dimension: viewDimension, baseMipLevel, mipLevelCount, - }); + }; - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); g.test('depth') @@ -184,11 +184,11 @@ fn getValue() -> u32 { view_type, mipCount ); - const view = texture.createView({ + const viewDescription = { dimension: viewDimension, baseMipLevel, mipLevelCount, - }); + }; - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, viewDescription, expected); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts index f1e9dd3793e0..24f7f9b6997f 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts @@ -50,9 +50,7 @@ fn getValue() -> u32 { `; const expected = [sampleCount]; - const view = texture.createView({}); - - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, {}, expected); }); g.test('depth') @@ -85,7 +83,5 @@ fn getValue() -> u32 { `; const expected = [sampleCount]; - const view = texture.createView({}); - - t.executeAndExpectResult(stage, code, view, expected); + t.executeAndExpectResult(stage, code, texture, {}, 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 354716cf6699..b6bf8949846c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -839,17 +839,19 @@ function getWeightForMipLevel( } /** - * Used for textureNumSamples, textureNumLevels, textureNumLayers + * Used for textureNumSamples, textureNumLevels, textureNumLayers, textureDimension */ export class WGSLTextureQueryTest extends GPUTest { executeAndExpectResult( stage: ShaderStage, code: string, - view: GPUTextureView, + texture: GPUTexture | GPUExternalTexture, + viewDescriptor: GPUTextureViewDescriptor | undefined, expected: number[] ) { const { device } = this; const returnType = `vec4`; + const castWGSL = `${returnType}(getValue()${range(4 - expected.length, () => ', 0').join('')})`; const stageWGSL = stage === 'vertex' ? ` @@ -860,7 +862,7 @@ export class WGSLTextureQueryTest extends GPUTest { let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, - ${returnType}(getValue())); + ${castWGSL}); } @fragment fn fsVertex(v: VOut) -> @location(0) vec4u { @@ -878,7 +880,7 @@ export class WGSLTextureQueryTest extends GPUTest { } @fragment fn fsFragment(v: VOut) -> @location(0) vec4u { - return bitcast(${returnType}(getValue())); + return bitcast(${castWGSL}); } ` : ` @@ -886,7 +888,7 @@ export class WGSLTextureQueryTest extends GPUTest { @group(1) @binding(0) var results: array<${returnType}>; @compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) { - results[id.x] = ${returnType}(getValue()); + results[id.x] = ${castWGSL}; } `; const wgsl = ` @@ -901,19 +903,96 @@ struct VOut { ${stageWGSL} `; const module = device.createShaderModule({ code: wgsl }); + + const visibility = + stage === 'compute' + ? GPUShaderStage.COMPUTE + : stage === 'fragment' + ? GPUShaderStage.FRAGMENT + : GPUShaderStage.VERTEX; + + const entries: GPUBindGroupLayoutEntry[] = []; + if (texture instanceof GPUExternalTexture) { + entries.push({ + binding: 0, + visibility, + externalTexture: {}, + }); + } else if (code.includes('texture_storage')) { + entries.push({ + binding: 0, + visibility, + storageTexture: { + access: code.includes(', read>') + ? 'read-only' + : code.includes(', write>') + ? 'write-only' + : 'read-write', + viewDimension: viewDescriptor?.dimension ?? '2d', + format: texture.format, + }, + }); + } else { + const sampleType = + viewDescriptor?.aspect === 'stencil-only' + ? 'uint' + : code.includes('texture_depth') + ? 'depth' + : isDepthTextureFormat(texture.format) + ? 'unfilterable-float' + : isStencilTextureFormat(texture.format) + ? 'uint' + : texture.sampleCount > 1 && kTextureFormatInfo[texture.format].color?.type === 'float' + ? 'unfilterable-float' + : kTextureFormatInfo[texture.format].color?.type ?? 'unfilterable-float'; + entries.push({ + binding: 0, + visibility, + texture: { + sampleType, + viewDimension: viewDescriptor?.dimension ?? '2d', + multisampled: texture.sampleCount > 1, + }, + }); + } + + const bindGroupLayouts: GPUBindGroupLayout[] = [device.createBindGroupLayout({ entries })]; + + if (stage === 'compute') { + bindGroupLayouts.push( + device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + hasDynamicOffset: false, + minBindingSize: 16, + }, + }, + ], + }) + ); + } + + const layout = device.createPipelineLayout({ + bindGroupLayouts, + }); + let pipeline: GPUComputePipeline | GPURenderPipeline; switch (stage) { case 'compute': pipeline = device.createComputePipeline({ - layout: 'auto', + layout, compute: { module }, }); break; case 'fragment': case 'vertex': pipeline = device.createRenderPipeline({ - layout: 'auto', + layout, vertex: { module }, fragment: { module, @@ -925,7 +1004,13 @@ struct VOut { const bindGroup0 = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), - entries: [{ binding: 0, resource: view }], + entries: [ + { + binding: 0, + resource: + texture instanceof GPUExternalTexture ? texture : texture.createView(viewDescriptor), + }, + ], }); const renderTarget = this.createTextureTracked({ @@ -990,7 +1075,7 @@ struct VOut { this.device.queue.submit([encoder.finish()]); const e = new Uint32Array(4); - e.set([expected[0], expected[0], expected[0], expected[0]]); + e.set(expected); this.expectGPUBufferValuesEqual(resultBuffer, e); } }