diff --git a/src/webgpu/api/operation/texture_view/write.spec.ts b/src/webgpu/api/operation/texture_view/write.spec.ts index 034012133479..b4ce6f4cec64 100644 --- a/src/webgpu/api/operation/texture_view/write.spec.ts +++ b/src/webgpu/api/operation/texture_view/write.spec.ts @@ -1,6 +1,9 @@ export const description = ` Test the result of writing textures through texture views with various options. +Reads value from a shader array, writes the value via various write methods. +Check the texture result with the expected texel view. + All x= every possible view write method: { - storage write {fragment, compute} - render pass store @@ -13,20 +16,358 @@ TODO: Write helper for this if not already available (see resource_init, buffer_ `; import { makeTestGroup } from '../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../gpu_test.js'; +import { unreachable } from '../../../../common/util/util.js'; +import { + kRegularTextureFormats, + kTextureFormatInfo, + RegularTextureFormat, +} from '../../../format_info.js'; +import { GPUTest, TextureTestMixin } from '../../../gpu_test.js'; +import { kFullscreenQuadVertexShaderCode } from '../../../util/shader.js'; +import { TexelView } from '../../../util/texture/texel_view.js'; + +export const g = makeTestGroup(TextureTestMixin(GPUTest)); + +const kTextureViewWriteMethods = [ + 'storage-write-fragment', + 'storage-write-compute', + 'render-pass-store', + 'render-pass-resolve', +] as const; +type TextureViewWriteMethod = (typeof kTextureViewWriteMethods)[number]; + +// Src color values to read from a shader array. +const kColorsFloat = [ + { R: 1.0, G: 0.0, B: 0.0, A: 0.8 }, + { R: 0.0, G: 1.0, B: 0.0, A: 0.7 }, + { R: 0.0, G: 0.0, B: 0.0, A: 0.6 }, + { R: 0.0, G: 0.0, B: 0.0, A: 0.5 }, + { R: 1.0, G: 1.0, B: 1.0, A: 0.4 }, + { R: 0.7, G: 0.0, B: 0.0, A: 0.3 }, + { R: 0.0, G: 0.8, B: 0.0, A: 0.2 }, + { R: 0.0, G: 0.0, B: 0.9, A: 0.1 }, + { R: 0.1, G: 0.2, B: 0.0, A: 0.3 }, + { R: 0.4, G: 0.3, B: 0.6, A: 0.8 }, +]; + +function FloatToIntColor(c: number) { + return Math.floor(c * 100); +} + +const kColorsInt = kColorsFloat.map(c => { + return { + R: FloatToIntColor(c.R), + G: FloatToIntColor(c.G), + B: FloatToIntColor(c.B), + A: FloatToIntColor(c.A), + }; +}); -export const g = makeTestGroup(GPUTest); +const kTextureSize = 16; + +function writeTextureAndGetExpectedTexelView( + t: GPUTest, + method: TextureViewWriteMethod, + view: GPUTextureView, + format: RegularTextureFormat, + sampleCount: number +) { + const info = kTextureFormatInfo[format]; + const isFloatType = info.color.type === 'float' || info.color.type === 'unfilterable-float'; + const kColors = isFloatType ? kColorsFloat : kColorsInt; + const expectedTexelView = TexelView.fromTexelsAsColors( + format, + coords => { + const pixelPos = coords.y * kTextureSize + coords.x; + return kColors[pixelPos % kColors.length]; + }, + { clampToFormatRange: true } + ); + const vecType = isFloatType ? 'vec4f' : info.color.type === 'sint' ? 'vec4i' : 'vec4u'; + const kColorArrayShaderString = `array<${vecType}, ${kColors.length}>( + ${kColors.map(t => `${vecType}(${t.R}, ${t.G}, ${t.B}, ${t.A}) `).join(',')} + )`; + + switch (method) { + case 'storage-write-compute': + { + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>; + @compute @workgroup_size(1, 1) fn main( + @builtin(global_invocation_id) global_id: vec3, + ) { + const src = ${kColorArrayShaderString}; + let coord = vec2u(global_id.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + textureStore(dst, coord, src[idx % ${kColors.length}]); + }`, + }), + entryPoint: 'main', + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup( + 0, + t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: view, + }, + ], + }) + ); + pass.dispatchWorkgroups(kTextureSize, kTextureSize); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + + case 'storage-write-fragment': + { + // Create a placeholder color attachment texture, + // The size of which equals that of format texture we are testing, + // so that we have the same number of fragments and texels. + const kPlaceholderTextureFormat = 'rgba8unorm'; + const placeholderTexture = t.trackForCleanup( + t.device.createTexture({ + format: kPlaceholderTextureFormat, + size: [kTextureSize, kTextureSize], + usage: GPUTextureUsage.RENDER_ATTACHMENT, + }) + ); + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: t.device.createShaderModule({ + code: kFullscreenQuadVertexShaderCode, + }), + }, + fragment: { + module: t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var dst: texture_storage_2d<${format}, write>; + @fragment fn main( + @builtin(position) fragCoord: vec4, + ) { + const src = ${kColorArrayShaderString}; + let coord = vec2u(fragCoord.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + textureStore(dst, coord, src[idx % ${kColors.length}]); + }`, + }), + // Set writeMask to 0 as the fragment shader has no output. + targets: [ + { + format: kPlaceholderTextureFormat, + writeMask: 0, + }, + ], + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: placeholderTexture.createView(), + loadOp: 'clear', + storeOp: 'discard', + }, + ], + }); + pass.setPipeline(pipeline); + pass.setBindGroup( + 0, + t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: view, + }, + ], + }) + ); + pass.draw(6); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + + case 'render-pass-store': + case 'render-pass-resolve': + { + // Create a placeholder color attachment texture for the store target when tesing texture is used as resolve target. + const targetView = + method === 'render-pass-store' + ? view + : t + .trackForCleanup( + t.device.createTexture({ + format, + size: [kTextureSize, kTextureSize], + usage: GPUTextureUsage.RENDER_ATTACHMENT, + sampleCount: 4, + }) + ) + .createView(); + const resolveView = method === 'render-pass-store' ? undefined : view; + const multisampleCount = method === 'render-pass-store' ? sampleCount : 4; + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: t.device.createShaderModule({ + code: kFullscreenQuadVertexShaderCode, + }), + }, + fragment: { + module: t.device.createShaderModule({ + code: ` + @fragment fn main( + @builtin(position) fragCoord: vec4, + ) -> @location(0) ${vecType} { + const src = ${kColorArrayShaderString}; + let coord = vec2u(fragCoord.xy); + let idx = coord.x + coord.y * ${kTextureSize}; + return src[idx % ${kColors.length}]; + }`, + }), + targets: [ + { + format, + }, + ], + }, + multisample: { + count: multisampleCount, + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: targetView, + resolveTarget: resolveView, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(pipeline); + pass.draw(6); + pass.end(); + t.device.queue.submit([commandEncoder.finish()]); + } + break; + default: + unreachable(); + } + + return expectedTexelView; +} g.test('format') .desc( `Views of every allowed format. +Read values from color array in the shader, and write it to the texture view via different write methods. + - x= every texture format - x= sampleCount {1, 4} if valid - x= every possible view write method (see above) + +TODO: Test sampleCount > 1 for 'render-pass-store' after extending copySinglePixelTextureToBufferUsingComputePass + to read multiple pixels from multisampled textures. [1] +TODO: Test rgb10a2uint when TexelRepresentation.numericRange is made per-component. [2] ` ) - .unimplemented(); + .params(u => + u // + .combine('method', kTextureViewWriteMethods) + .combine('format', kRegularTextureFormats) + .combine('sampleCount', [1, 4]) + .filter(({ format, method, sampleCount }) => { + const info = kTextureFormatInfo[format]; + + if (sampleCount > 1 && !info.multisample) { + return false; + } + + // [2] + if (format === 'rgb10a2uint') { + return false; + } + + switch (method) { + case 'storage-write-compute': + case 'storage-write-fragment': + return info.color?.storage && sampleCount === 1; + case 'render-pass-store': + // [1] + if (sampleCount > 1) { + return false; + } + return !!info.colorRender; + case 'render-pass-resolve': + return !!info.colorRender?.resolve && sampleCount === 1; + } + return true; + }) + ) + .beforeAllSubcases(t => { + const { format, method } = t.params; + t.skipIfTextureFormatNotSupported(format); + + switch (method) { + case 'storage-write-compute': + case 'storage-write-fragment': + // Still need to filter again for compat mode. + t.skipIfTextureFormatNotUsableAsStorageTexture(format); + break; + } + }) + .fn(t => { + const { format, method, sampleCount } = t.params; + + const usage = + GPUTextureUsage.COPY_SRC | + (method.includes('storage') + ? GPUTextureUsage.STORAGE_BINDING + : GPUTextureUsage.RENDER_ATTACHMENT); + + const texture = t.trackForCleanup( + t.device.createTexture({ + format, + usage, + size: [kTextureSize, kTextureSize], + sampleCount, + }) + ); + + const view = texture.createView(); + const expectedTexelView = writeTextureAndGetExpectedTexelView( + t, + method, + view, + format, + sampleCount + ); + + // [1] Use copySinglePixelTextureToBufferUsingComputePass to check multisampled texture. + t.expectTexelViewComparisonIsOkInTexture({ texture }, expectedTexelView, [ + kTextureSize, + kTextureSize, + ]); + }); g.test('dimension') .desc( diff --git a/src/webgpu/util/shader.ts b/src/webgpu/util/shader.ts index 2a09061527c5..6f291ec4d1c3 100644 --- a/src/webgpu/util/shader.ts +++ b/src/webgpu/util/shader.ts @@ -11,6 +11,27 @@ export const kDefaultFragmentShaderCode = ` return vec4(1.0, 1.0, 1.0, 1.0); }`; +// MAINTENANCE_TODO(#3344): deduplicate fullscreen quad shader code. +export const kFullscreenQuadVertexShaderCode = ` + struct VertexOutput { + @builtin(position) Position : vec4 + }; + + @vertex fn main(@builtin(vertex_index) VertexIndex : u32) -> VertexOutput { + var pos = array, 6>( + vec2( 1.0, 1.0), + vec2( 1.0, -1.0), + vec2(-1.0, -1.0), + vec2( 1.0, 1.0), + vec2(-1.0, -1.0), + vec2(-1.0, 1.0)); + + var output : VertexOutput; + output.Position = vec4(pos[VertexIndex], 0.0, 1.0); + return output; + } +`; + const kPlainTypeInfo = { i32: { suffix: '',