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 500376321444..dd5d5923172f 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLayers.spec.ts @@ -7,6 +7,7 @@ Returns the number of layers (elements) of an array texture. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kTextureFormatInfo } from '../../../../../format_info.js'; import { TexelFormats } from '../../../../types.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -54,6 +55,7 @@ Parameters .combine('view_type', ['full', 'partial'] as const) .beginSubcases() .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIf( @@ -66,7 +68,7 @@ Parameters ); }) .fn(t => { - const { texture_type, sampled_type, view_type } = t.params; + const { stage, texture_type, sampled_type, view_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const texture = t.createTextureTracked({ @@ -77,9 +79,8 @@ Parameters 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); +fn getValue() -> u32 { + return textureNumLayers(t); } `; @@ -93,7 +94,7 @@ Parameters arrayLayerCount, }); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, expected); }); g.test('arrayed') @@ -111,6 +112,8 @@ Parameters u .combine('texture_type', ['texture_depth_2d_array', 'texture_depth_cube_array'] as const) .combine('view_type', ['full', 'partial'] as const) + .beginSubcases() + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIf( @@ -123,7 +126,7 @@ Parameters ); }) .fn(t => { - const { texture_type, view_type } = t.params; + const { stage, texture_type, view_type } = t.params; const texture = t.createTextureTracked({ format: 'depth32float', @@ -134,8 +137,8 @@ Parameters 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); +fn getValue() -> u32 { + return textureNumLayers(t); } `; @@ -149,7 +152,7 @@ Parameters arrayLayerCount, }); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, expected); }); g.test('storage') @@ -185,10 +188,13 @@ Parameters .combineWithParams(TexelFormats) .combine('view_type', ['full', 'partial'] as const) .beginSubcases() + .combine('stage', kShaderStages) .combine('access_mode', ['read', 'write', 'read_write'] as const) .filter( t => t.access_mode !== 'read_write' || kTextureFormatInfo[t.format].color?.readWriteStorage ) + // Vertex stage can not use writable storage textures. + .unless(t => t.stage === 'vertex' && t.access_mode !== 'read') ) .beforeAllSubcases(t => { t.skipIf( @@ -198,7 +204,7 @@ Parameters t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format); }) .fn(t => { - const { format, access_mode, view_type } = t.params; + const { stage, format, access_mode, view_type } = t.params; const texture = t.createTextureTracked({ format, @@ -209,8 +215,8 @@ Parameters 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); +fn getValue() -> u32 { + return textureNumLayers(t); } `; @@ -223,5 +229,5 @@ Parameters arrayLayerCount, }); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, 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 471a462504d4..61fb1f745646 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumLevels.spec.ts @@ -6,6 +6,7 @@ Returns the number of mip levels of a texture. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { getTextureDimensionFromView } from '../../../../../util/texture/base.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -66,6 +67,7 @@ Parameters 'texture_cube_array', ] as const) .beginSubcases() + .combine('stage', kShaderStages) .combine('sampled_type', ['f32', 'i32', 'u32'] as const) .combine('view_type', ['full', 'partial'] as const) // 1d textures can't have mipLevelCount > 0 @@ -75,7 +77,7 @@ Parameters t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); }) .fn(t => { - const { texture_type, sampled_type, view_type } = t.params; + const { stage, texture_type, sampled_type, view_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const viewDimension = kTextureTypeToViewDimension[texture_type]; @@ -101,8 +103,8 @@ Parameters 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 = textureNumLevels(t); +fn getValue() -> u32 { + return textureNumLevels(t); } `; @@ -116,7 +118,7 @@ Parameters mipLevelCount, }); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, expected); }); g.test('depth') @@ -141,12 +143,14 @@ Parameters 'texture_depth_cube_array', ] as const) .combine('view_type', ['full', 'partial'] as const) + .beginSubcases() + .combine('stage', kShaderStages) ) .beforeAllSubcases(t => { t.skipIfTextureViewDimensionNotSupported(kTextureTypeToViewDimension[t.params.texture_type]); }) .fn(t => { - const { texture_type, view_type } = t.params; + const { stage, texture_type, view_type } = t.params; const viewDimension = kTextureTypeToViewDimension[texture_type]; const dimension = getTextureDimensionFromView(viewDimension); @@ -171,8 +175,8 @@ Parameters const code = ` @group(0) @binding(0) var t: ${texture_type}; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumLevels(t); +fn getValue() -> u32 { + return textureNumLevels(t); } `; @@ -186,5 +190,5 @@ Parameters mipLevelCount, }); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, 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 a6314198529b..f1e9dd3793e0 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureNumSamples.spec.ts @@ -5,6 +5,7 @@ Returns the number samples per texel in a multisampled texture. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { kSampleTypeInfo, WGSLTextureQueryTest } from './texture_utils.js'; @@ -22,9 +23,14 @@ Parameters * t The multisampled texture. ` ) - .params(u => u.beginSubcases().combine('sampled_type', ['f32', 'i32', 'u32'] as const)) + .params(u => + u + .beginSubcases() + .combine('stage', kShaderStages) + .combine('sampled_type', ['f32', 'i32', 'u32'] as const) + ) .fn(t => { - const { sampled_type } = t.params; + const { stage, sampled_type } = t.params; const { format } = kSampleTypeInfo[sampled_type]; const sampleCount = 4; @@ -38,15 +44,15 @@ Parameters const code = ` @group(0) @binding(0) var t: texture_multisampled_2d<${sampled_type}>; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumSamples(t); +fn getValue() -> u32 { + return textureNumSamples(t); } `; const expected = [sampleCount]; const view = texture.createView({}); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, code, view, expected); }); g.test('depth') @@ -59,7 +65,9 @@ Parameters * t The multisampled texture. ` ) + .params(u => u.beginSubcases().combine('stage', kShaderStages)) .fn(t => { + const { stage } = t.params; const sampleCount = 4; const texture = t.createTextureTracked({ format: 'depth32float', @@ -71,13 +79,13 @@ Parameters const code = ` @group(0) @binding(0) var t: texture_depth_multisampled_2d; @group(0) @binding(1) var result: u32; -@compute @workgroup_size(1) fn cs() { - result = textureNumSamples(t); +fn getValue() -> u32 { + return textureNumSamples(t); } `; const expected = [sampleCount]; const view = texture.createView({}); - t.executeAndExpectResult(code, view, expected); + t.executeAndExpectResult(stage, 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 90d2f937126c..354716cf6699 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -839,42 +839,158 @@ function getWeightForMipLevel( } /** - * Used for textureDimension, textureNumLevels, textureNumLayers + * Used for textureNumSamples, textureNumLevels, textureNumLayers */ export class WGSLTextureQueryTest extends GPUTest { - executeAndExpectResult(code: string, view: GPUTextureView, expected: number[]) { + executeAndExpectResult( + stage: ShaderStage, + code: string, + view: GPUTextureView, + expected: number[] + ) { const { device } = this; - const module = device.createShaderModule({ code }); - const pipeline = device.createComputePipeline({ - layout: 'auto', - compute: { - module, - }, + const returnType = `vec4`; + const stageWGSL = + stage === 'vertex' + ? ` +// --------------------------- vertex stage shaders -------------------------------- +@vertex fn vsVertex( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), + instance_index, + ${returnType}(getValue())); +} + +@fragment fn fsVertex(v: VOut) -> @location(0) vec4u { + return bitcast(v.result); +} +` + : stage === 'fragment' + ? ` +// --------------------------- fragment stage shaders -------------------------------- +@vertex fn vsFragment( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0)); +} + +@fragment fn fsFragment(v: VOut) -> @location(0) vec4u { + return bitcast(${returnType}(getValue())); +} +` + : ` +// --------------------------- compute stage shaders -------------------------------- +@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()); +} +`; + const wgsl = ` + ${code} + +struct VOut { + @builtin(position) pos: vec4f, + @location(0) @interpolate(flat, either) ndx: u32, + @location(1) @interpolate(flat, either) result: ${returnType}, +}; + + ${stageWGSL} + `; + const module = device.createShaderModule({ code: wgsl }); + let pipeline: GPUComputePipeline | GPURenderPipeline; + + switch (stage) { + case 'compute': + pipeline = device.createComputePipeline({ + layout: 'auto', + compute: { module }, + }); + break; + case 'fragment': + case 'vertex': + pipeline = device.createRenderPipeline({ + layout: 'auto', + vertex: { module }, + fragment: { + module, + targets: [{ format: 'rgba32uint' }], + }, + }); + break; + } + + const bindGroup0 = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [{ binding: 0, resource: view }], }); - const resultBuffer = this.createBufferTracked({ - size: 16, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + const renderTarget = this.createTextureTracked({ + format: 'rgba32uint', + size: [expected.length, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, }); - const bindGroup = device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { binding: 0, resource: view }, - { binding: 1, resource: { buffer: resultBuffer } }, - ], + const resultBuffer = this.createBufferTracked({ + size: align(expected.length * 4, 256), + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, }); + let storageBuffer: GPUBuffer | undefined; 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()]); + + if (stage === 'compute') { + storageBuffer = this.createBufferTracked({ + size: resultBuffer.size, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bindGroup1 = device.createBindGroup({ + layout: pipeline!.getBindGroupLayout(1), + entries: [{ binding: 0, resource: { buffer: storageBuffer } }], + }); + + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline! as GPUComputePipeline); + pass.setBindGroup(0, bindGroup0); + pass.setBindGroup(1, bindGroup1); + pass.dispatchWorkgroups(expected.length); + pass.end(); + encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size); + } else { + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTarget.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + + pass.setPipeline(pipeline! as GPURenderPipeline); + pass.setBindGroup(0, bindGroup0); + for (let i = 0; i < expected.length; ++i) { + pass.setViewport(i, 0, 1, 1, 0, 1); + pass.draw(3, 1, 0, i); + } + pass.end(); + encoder.copyTextureToBuffer( + { texture: renderTarget }, + { + buffer: resultBuffer, + bytesPerRow: resultBuffer.size, + }, + [renderTarget.width, 1] + ); + } + this.device.queue.submit([encoder.finish()]); const e = new Uint32Array(4); - e.set(expected); + e.set([expected[0], expected[0], expected[0], expected[0]]); this.expectGPUBufferValuesEqual(resultBuffer, e); } }