diff options
Diffstat (limited to '')
5 files changed, 1964 insertions, 0 deletions
diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/buffer.spec.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/buffer.spec.ts new file mode 100644 index 0000000000..73c50b8393 --- /dev/null +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/buffer.spec.ts @@ -0,0 +1,899 @@ +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { unreachable } from '../../../../common/util/util.js'; +import { GPUConst } from '../../../constants.js'; +import { GPUTest } from '../../../gpu_test.js'; +import { getTextureCopyLayout } from '../../../util/texture/layout.js'; +import { PerTexelComponent } from '../../../util/texture/texel_data.js'; + +export const description = ` +Test uninitialized buffers are initialized to zero when read +(or read-written, e.g. with depth write or atomics). + +Note that: +- We don't need 'copy_buffer_to_buffer_copy_destination' here because there has already been an + operation test 'command_buffer.copyBufferToBuffer.single' that provides the same functionality. +`; + +const kMapModeOptions = [GPUConst.MapMode.READ, GPUConst.MapMode.WRITE]; +const kBufferUsagesForMappedAtCreationTests = [ + GPUConst.BufferUsage.COPY_DST | GPUConst.BufferUsage.MAP_READ, + GPUConst.BufferUsage.COPY_SRC | GPUConst.BufferUsage.MAP_WRITE, + GPUConst.BufferUsage.COPY_SRC, +]; + +class F extends GPUTest { + GetBufferUsageFromMapMode(mapMode: GPUMapModeFlags): number { + switch (mapMode) { + case GPUMapMode.READ: + return GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ; + case GPUMapMode.WRITE: + return GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE; + default: + unreachable(); + return 0; + } + } + + async CheckGPUBufferContent( + buffer: GPUBuffer, + bufferUsage: GPUBufferUsageFlags, + expectedData: Uint8Array + ): Promise<void> { + const mappable = bufferUsage & GPUBufferUsage.MAP_READ; + this.expectGPUBufferValuesEqual(buffer, expectedData, 0, { method: mappable ? 'map' : 'copy' }); + } + + TestBufferZeroInitInBindGroup( + computeShaderModule: GPUShaderModule, + buffer: GPUBuffer, + bufferOffset: number, + boundBufferSize: number + ): void { + const computePipeline = this.device.createComputePipeline({ + layout: 'auto', + compute: { + module: computeShaderModule, + entryPoint: 'main', + }, + }); + const outputTexture = this.device.createTexture({ + format: 'rgba8unorm', + size: [1, 1, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.STORAGE_BINDING, + }); + this.trackForCleanup(outputTexture); + const bindGroup = this.device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer, + offset: bufferOffset, + size: boundBufferSize, + }, + }, + { + binding: 1, + resource: outputTexture.createView(), + }, + ], + }); + + const encoder = this.device.createCommandEncoder(); + const computePass = encoder.beginComputePass(); + computePass.setBindGroup(0, bindGroup); + computePass.setPipeline(computePipeline); + computePass.dispatchWorkgroups(1); + computePass.end(); + this.queue.submit([encoder.finish()]); + + this.CheckBufferAndOutputTexture(buffer, boundBufferSize + bufferOffset, outputTexture); + } + + CreateRenderPipelineForTest( + vertexShaderModule: GPUShaderModule, + testVertexBuffer: boolean + ): GPURenderPipeline { + const renderPipelineDescriptor: GPURenderPipelineDescriptor = { + layout: 'auto', + vertex: { + module: vertexShaderModule, + entryPoint: 'main', + }, + fragment: { + module: this.device.createShaderModule({ + code: ` + @fragment + fn main(@location(0) i_color : vec4<f32>) -> @location(0) vec4<f32> { + return i_color; + }`, + }), + entryPoint: 'main', + targets: [{ format: 'rgba8unorm' }], + }, + primitive: { + topology: 'point-list', + }, + }; + if (testVertexBuffer) { + renderPipelineDescriptor.vertex.buffers = [ + { + arrayStride: 16, + attributes: [{ format: 'float32x4', offset: 0, shaderLocation: 0 }], + }, + ]; + } + + return this.device.createRenderPipeline(renderPipelineDescriptor); + } + + RecordInitializeTextureColor( + encoder: GPUCommandEncoder, + texture: GPUTexture, + color: GPUColor + ): void { + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: texture.createView(), + clearValue: color, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + renderPass.end(); + } + + CheckBufferAndOutputTexture( + buffer: GPUBuffer, + bufferSize: number, + outputTexture: GPUTexture, + outputTextureSize: [number, number, number] = [1, 1, 1], + outputTextureColor: PerTexelComponent<number> = { R: 0.0, G: 1.0, B: 0.0, A: 1.0 } + ): void { + this.expectSingleColor(outputTexture, 'rgba8unorm', { + size: outputTextureSize, + exp: outputTextureColor, + }); + + const expectedBufferData = new Uint8Array(bufferSize); + this.expectGPUBufferValuesEqual(buffer, expectedBufferData); + } +} + +export const g = makeTestGroup(F); + +g.test('partial_write_buffer') + .desc( + `Verify when we upload data to a part of a buffer with writeBuffer() just after the creation of +the buffer, the remaining part of that buffer will be initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('offset', [0, 8, -12])) + .fn(async t => { + const { offset } = t.params; + const bufferSize = 32; + const appliedOffset = offset >= 0 ? offset : bufferSize + offset; + + const buffer = t.device.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + }); + t.trackForCleanup(buffer); + + const copySize = 12; + const writeData = new Uint8Array(copySize); + const expectedData = new Uint8Array(bufferSize); + for (let i = 0; i < copySize; ++i) { + expectedData[appliedOffset + i] = writeData[i] = i + 1; + } + t.queue.writeBuffer(buffer, appliedOffset, writeData, 0); + + t.expectGPUBufferValuesEqual(buffer, expectedData); + }); + +g.test('map_whole_buffer') + .desc( + `Verify when we map the whole range of a mappable GPUBuffer to a typed array buffer just after +creating the GPUBuffer, the contents of both the typed array buffer and the GPUBuffer itself +have already been initialized to 0.` + ) + .params(u => u.combine('mapMode', kMapModeOptions)) + .fn(async t => { + const { mapMode } = t.params; + + const bufferSize = 32; + const bufferUsage = t.GetBufferUsageFromMapMode(mapMode); + const buffer = t.device.createBuffer({ + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(buffer); + + await buffer.mapAsync(mapMode); + const readData = new Uint8Array(buffer.getMappedRange()); + for (let i = 0; i < bufferSize; ++i) { + t.expect(readData[i] === 0); + } + buffer.unmap(); + + const expectedData = new Uint8Array(bufferSize); + await t.CheckGPUBufferContent(buffer, bufferUsage, expectedData); + }); + +g.test('map_partial_buffer') + .desc( + `Verify when we map a subrange of a mappable GPUBuffer to a typed array buffer just after the +creation of the GPUBuffer, the contents of both the typed array buffer and the GPUBuffer have +already been initialized to 0.` + ) + .params(u => u.combine('mapMode', kMapModeOptions).beginSubcases().combine('offset', [0, 8, -16])) + .fn(async t => { + const { mapMode, offset } = t.params; + const bufferSize = 32; + const appliedOffset = offset >= 0 ? offset : bufferSize + offset; + + const bufferUsage = t.GetBufferUsageFromMapMode(mapMode); + const buffer = t.device.createBuffer({ + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(buffer); + + const expectedData = new Uint8Array(bufferSize); + { + const mapSize = 16; + await buffer.mapAsync(mapMode, appliedOffset, mapSize); + const mappedData = new Uint8Array(buffer.getMappedRange(appliedOffset, mapSize)); + for (let i = 0; i < mapSize; ++i) { + t.expect(mappedData[i] === 0); + if (mapMode === GPUMapMode.WRITE) { + mappedData[i] = expectedData[appliedOffset + i] = i + 1; + } + } + buffer.unmap(); + } + + await t.CheckGPUBufferContent(buffer, bufferUsage, expectedData); + }); + +g.test('mapped_at_creation_whole_buffer') + .desc( + `Verify when we call getMappedRange() at the whole range of a GPUBuffer created with +mappedAtCreation === true just after its creation, the contents of both the returned typed +array buffer of getMappedRange() and the GPUBuffer itself have all been initialized to 0.` + ) + .params(u => u.combine('bufferUsage', kBufferUsagesForMappedAtCreationTests)) + .fn(async t => { + const { bufferUsage } = t.params; + + const bufferSize = 32; + const buffer = t.device.createBuffer({ + mappedAtCreation: true, + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(buffer); + + const mapped = new Uint8Array(buffer.getMappedRange()); + for (let i = 0; i < bufferSize; ++i) { + t.expect(mapped[i] === 0); + } + buffer.unmap(); + + const expectedData = new Uint8Array(bufferSize); + await t.CheckGPUBufferContent(buffer, bufferUsage, expectedData); + }); + +g.test('mapped_at_creation_partial_buffer') + .desc( + `Verify when we call getMappedRange() at a subrange of a GPUBuffer created with +mappedAtCreation === true just after its creation, the contents of both the returned typed +array buffer of getMappedRange() and the GPUBuffer itself have all been initialized to 0.` + ) + .params(u => + u + .combine('bufferUsage', kBufferUsagesForMappedAtCreationTests) + .beginSubcases() + .combine('offset', [0, 8, -16]) + ) + .fn(async t => { + const { bufferUsage, offset } = t.params; + const bufferSize = 32; + const appliedOffset = offset >= 0 ? offset : bufferSize + offset; + + const buffer = t.device.createBuffer({ + mappedAtCreation: true, + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(buffer); + + const expectedData = new Uint8Array(bufferSize); + { + const mappedSize = 12; + const mapped = new Uint8Array(buffer.getMappedRange(appliedOffset, mappedSize)); + for (let i = 0; i < mappedSize; ++i) { + t.expect(mapped[i] === 0); + if (!(bufferUsage & GPUBufferUsage.MAP_READ)) { + mapped[i] = expectedData[appliedOffset + i] = i + 1; + } + } + buffer.unmap(); + } + + await t.CheckGPUBufferContent(buffer, bufferUsage, expectedData); + }); + +g.test('copy_buffer_to_buffer_copy_source') + .desc( + `Verify when the first usage of a GPUBuffer is being used as the source buffer of +CopyBufferToBuffer(), the contents of the GPUBuffer have already been initialized to 0.` + ) + .fn(async t => { + const bufferSize = 32; + const bufferUsage = GPUBufferUsage.COPY_SRC; + const buffer = t.device.createBuffer({ + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(buffer); + + const expectedData = new Uint8Array(bufferSize); + // copyBufferToBuffer() is called inside t.CheckGPUBufferContent(). + await t.CheckGPUBufferContent(buffer, bufferUsage, expectedData); + }); + +g.test('copy_buffer_to_texture') + .desc( + `Verify when the first usage of a GPUBuffer is being used as the source buffer of +CopyBufferToTexture(), the contents of the GPUBuffer have already been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 8])) + .fn(async t => { + const { bufferOffset } = t.params; + const textureSize: [number, number, number] = [8, 8, 1]; + const dstTextureFormat = 'rgba8unorm'; + + const dstTexture = t.device.createTexture({ + size: textureSize, + format: dstTextureFormat, + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST, + }); + t.trackForCleanup(dstTexture); + const layout = getTextureCopyLayout(dstTextureFormat, '2d', textureSize); + const srcBufferSize = layout.byteLength + bufferOffset; + const srcBufferUsage = GPUBufferUsage.COPY_SRC; + const srcBuffer = t.device.createBuffer({ + size: srcBufferSize, + usage: srcBufferUsage, + }); + t.trackForCleanup(srcBuffer); + + const encoder = t.device.createCommandEncoder(); + encoder.copyBufferToTexture( + { + buffer: srcBuffer, + offset: bufferOffset, + bytesPerRow: layout.bytesPerRow, + rowsPerImage: layout.rowsPerImage, + }, + { texture: dstTexture }, + textureSize + ); + t.queue.submit([encoder.finish()]); + + t.CheckBufferAndOutputTexture(srcBuffer, srcBufferSize, dstTexture, textureSize, { + R: 0.0, + G: 0.0, + B: 0.0, + A: 0.0, + }); + }); + +g.test('resolve_query_set_to_partial_buffer') + .desc( + `Verify when we resolve a query set into a GPUBuffer just after creating that GPUBuffer, the +remaining part of it will be initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 256])) + .fn(async t => { + const { bufferOffset } = t.params; + const bufferSize = bufferOffset + 8; + const bufferUsage = GPUBufferUsage.COPY_SRC | GPUBufferUsage.QUERY_RESOLVE; + const dstBuffer = t.device.createBuffer({ + size: bufferSize, + usage: bufferUsage, + }); + t.trackForCleanup(dstBuffer); + + const querySet = t.device.createQuerySet({ type: 'occlusion', count: 1 }); + const encoder = t.device.createCommandEncoder(); + encoder.resolveQuerySet(querySet, 0, 1, dstBuffer, bufferOffset); + t.queue.submit([encoder.finish()]); + + const expectedBufferData = new Uint8Array(bufferSize); + await t.CheckGPUBufferContent(dstBuffer, bufferUsage, expectedBufferData); + }); + +g.test('copy_texture_to_partial_buffer') + .desc( + `Verify when we copy from a GPUTexture into a GPUBuffer just after creating that GPUBuffer, the +remaining part of it will be initialized to 0.` + ) + .paramsSubcasesOnly(u => + u + .combine('bufferOffset', [0, 8, -16]) + .combine('arrayLayerCount', [1, 3]) + .combine('copyMipLevel', [0, 2]) + .combine('rowsPerImage', [16, 20]) + .filter(t => { + // We don't need to test the copies that will cover the whole GPUBuffer. + return !(t.bufferOffset === 0 && t.rowsPerImage === 16); + }) + ) + .fn(async t => { + const { bufferOffset, arrayLayerCount, copyMipLevel, rowsPerImage } = t.params; + const srcTextureFormat = 'r8uint'; + const textureSize = [32, 16, arrayLayerCount] as const; + + const srcTexture = t.device.createTexture({ + format: srcTextureFormat, + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, + size: textureSize, + mipLevelCount: copyMipLevel + 1, + }); + t.trackForCleanup(srcTexture); + + const bytesPerRow = 256; + const layout = getTextureCopyLayout(srcTextureFormat, '2d', textureSize, { + mipLevel: copyMipLevel, + bytesPerRow, + rowsPerImage, + }); + + const dstBufferSize = layout.byteLength + Math.abs(bufferOffset); + const dstBuffer = t.device.createBuffer({ + size: dstBufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST, + }); + t.trackForCleanup(dstBuffer); + + const encoder = t.device.createCommandEncoder(); + + // Initialize srcTexture + for (let layer = 0; layer < arrayLayerCount; ++layer) { + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: srcTexture.createView({ + baseArrayLayer: layer, + arrayLayerCount: 1, + baseMipLevel: copyMipLevel, + }), + clearValue: { r: layer + 1, g: 0, b: 0, a: 0 }, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + renderPass.end(); + } + + // Do texture-to-buffer copy + const appliedOffset = Math.max(bufferOffset, 0); + encoder.copyTextureToBuffer( + { texture: srcTexture, mipLevel: copyMipLevel }, + { buffer: dstBuffer, offset: appliedOffset, bytesPerRow, rowsPerImage }, + layout.mipSize + ); + t.queue.submit([encoder.finish()]); + + // Check if the contents of the destination buffer are what we expect. + const expectedData = new Uint8Array(dstBufferSize); + for (let layer = 0; layer < arrayLayerCount; ++layer) { + for (let y = 0; y < layout.mipSize[1]; ++y) { + for (let x = 0; x < layout.mipSize[0]; ++x) { + expectedData[appliedOffset + layer * bytesPerRow * rowsPerImage + y * bytesPerRow + x] = + layer + 1; + } + } + } + t.expectGPUBufferValuesEqual(dstBuffer, expectedData); + }); + +g.test('uniform_buffer') + .desc( + `Verify when we use a GPUBuffer as a uniform buffer just after the creation of that GPUBuffer, + all the contents in that GPUBuffer have been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 256])) + .fn(async t => { + const { bufferOffset } = t.params; + + const boundBufferSize = 16; + const buffer = t.device.createBuffer({ + size: bufferOffset + boundBufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.UNIFORM, + }); + t.trackForCleanup(buffer); + + const computeShaderModule = t.device.createShaderModule({ + code: ` + struct UBO { + value : vec4<u32> + }; + @group(0) @binding(0) var<uniform> ubo : UBO; + @group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>; + + @compute @workgroup_size(1) fn main() { + if (all(ubo.value == vec4<u32>(0u, 0u, 0u, 0u))) { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0)); + } else { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(1.0, 0.0, 0.0, 1.0)); + } + }`, + }); + + // Verify the whole range of the buffer has been initialized to 0 in a compute shader. + t.TestBufferZeroInitInBindGroup(computeShaderModule, buffer, bufferOffset, boundBufferSize); + }); + +g.test('readonly_storage_buffer') + .desc( + `Verify when we use a GPUBuffer as a read-only storage buffer just after the creation of that + GPUBuffer, all the contents in that GPUBuffer have been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 256])) + .fn(async t => { + const { bufferOffset } = t.params; + const boundBufferSize = 16; + const buffer = t.device.createBuffer({ + size: bufferOffset + boundBufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, + }); + t.trackForCleanup(buffer); + + const computeShaderModule = t.device.createShaderModule({ + code: ` + struct SSBO { + value : vec4<u32> + }; + @group(0) @binding(0) var<storage, read> ssbo : SSBO; + @group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>; + + @compute @workgroup_size(1) fn main() { + if (all(ssbo.value == vec4<u32>(0u, 0u, 0u, 0u))) { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0)); + } else { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(1.0, 0.0, 0.0, 1.0)); + } + }`, + }); + + // Verify the whole range of the buffer has been initialized to 0 in a compute shader. + t.TestBufferZeroInitInBindGroup(computeShaderModule, buffer, bufferOffset, boundBufferSize); + }); + +g.test('storage_buffer') + .desc( + `Verify when we use a GPUBuffer as a storage buffer just after the creation of that + GPUBuffer, all the contents in that GPUBuffer have been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 256])) + .fn(async t => { + const { bufferOffset } = t.params; + const boundBufferSize = 16; + const buffer = t.device.createBuffer({ + size: bufferOffset + boundBufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, + }); + t.trackForCleanup(buffer); + + const computeShaderModule = t.device.createShaderModule({ + code: ` + struct SSBO { + value : vec4<u32> + }; + @group(0) @binding(0) var<storage, read_write> ssbo : SSBO; + @group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>; + + @compute @workgroup_size(1) fn main() { + if (all(ssbo.value == vec4<u32>(0u, 0u, 0u, 0u))) { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0)); + } else { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(1.0, 0.0, 0.0, 1.0)); + } + }`, + }); + + // Verify the whole range of the buffer has been initialized to 0 in a compute shader. + t.TestBufferZeroInitInBindGroup(computeShaderModule, buffer, bufferOffset, boundBufferSize); + }); + +g.test('vertex_buffer') + .desc( + `Verify when we use a GPUBuffer as a vertex buffer just after the creation of that + GPUBuffer, all the contents in that GPUBuffer have been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 16])) + .fn(async t => { + const { bufferOffset } = t.params; + + const renderPipeline = t.CreateRenderPipelineForTest( + t.device.createShaderModule({ + code: ` + struct VertexOut { + @location(0) color : vec4<f32>, + @builtin(position) position : vec4<f32>, + }; + + @vertex fn main(@location(0) pos : vec4<f32>) -> VertexOut { + var output : VertexOut; + if (all(pos == vec4<f32>(0.0, 0.0, 0.0, 0.0))) { + output.color = vec4<f32>(0.0, 1.0, 0.0, 1.0); + } else { + output.color = vec4<f32>(1.0, 0.0, 0.0, 1.0); + } + output.position = vec4<f32>(0.0, 0.0, 0.0, 1.0); + return output; + }`, + }), + true + ); + + const bufferSize = 16 + bufferOffset; + const vertexBuffer = t.device.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_SRC, + }); + t.trackForCleanup(vertexBuffer); + + const outputTexture = t.device.createTexture({ + format: 'rgba8unorm', + size: [1, 1, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, + }); + t.trackForCleanup(outputTexture); + + const encoder = t.device.createCommandEncoder(); + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: outputTexture.createView(), + clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 0.0 }, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + renderPass.setVertexBuffer(0, vertexBuffer, bufferOffset); + renderPass.setPipeline(renderPipeline); + renderPass.draw(1); + renderPass.end(); + t.queue.submit([encoder.finish()]); + + t.CheckBufferAndOutputTexture(vertexBuffer, bufferSize, outputTexture); + }); + +g.test('index_buffer') + .desc( + `Verify when we use a GPUBuffer as an index buffer just after the creation of that +GPUBuffer, all the contents in that GPUBuffer have been initialized to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 16])) + .fn(async t => { + const { bufferOffset } = t.params; + + const renderPipeline = t.CreateRenderPipelineForTest( + t.device.createShaderModule({ + code: ` + struct VertexOut { + @location(0) color : vec4<f32>, + @builtin(position) position : vec4<f32>, + }; + + @vertex + fn main(@builtin(vertex_index) VertexIndex : u32) -> VertexOut { + var output : VertexOut; + if (VertexIndex == 0u) { + output.color = vec4<f32>(0.0, 1.0, 0.0, 1.0); + } else { + output.color = vec4<f32>(1.0, 0.0, 0.0, 1.0); + } + output.position = vec4<f32>(0.0, 0.0, 0.0, 1.0); + return output; + }`, + }), + false + ); + + // The size of GPUBuffer must be at least 4. + const bufferSize = 4 + bufferOffset; + const indexBuffer = t.device.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.INDEX | GPUBufferUsage.COPY_SRC, + }); + t.trackForCleanup(indexBuffer); + + const outputTexture = t.device.createTexture({ + format: 'rgba8unorm', + size: [1, 1, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, + }); + t.trackForCleanup(outputTexture); + + const encoder = t.device.createCommandEncoder(); + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: outputTexture.createView(), + clearValue: { r: 0.0, g: 0.0, b: 0.0, a: 0.0 }, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + renderPass.setPipeline(renderPipeline); + renderPass.setIndexBuffer(indexBuffer, 'uint16', bufferOffset, 4); + renderPass.drawIndexed(1); + renderPass.end(); + t.queue.submit([encoder.finish()]); + + t.CheckBufferAndOutputTexture(indexBuffer, bufferSize, outputTexture); + }); + +g.test('indirect_buffer_for_draw_indirect') + .desc( + `Verify when we use a GPUBuffer as an indirect buffer for drawIndirect() or +drawIndexedIndirect() just after the creation of that GPUBuffer, all the contents in that GPUBuffer +have been initialized to 0.` + ) + .params(u => + u.combine('test_indexed_draw', [true, false]).beginSubcases().combine('bufferOffset', [0, 16]) + ) + .fn(async t => { + const { test_indexed_draw, bufferOffset } = t.params; + + const renderPipeline = t.CreateRenderPipelineForTest( + t.device.createShaderModule({ + code: ` + struct VertexOut { + @location(0) color : vec4<f32>, + @builtin(position) position : vec4<f32>, + }; + + @vertex fn main() -> VertexOut { + var output : VertexOut; + output.color = vec4<f32>(1.0, 0.0, 0.0, 1.0); + output.position = vec4<f32>(0.0, 0.0, 0.0, 1.0); + return output; + }`, + }), + false + ); + + const kDrawIndirectParametersSize = 16; + const kDrawIndexedIndirectParametersSize = 20; + const bufferSize = + Math.max(kDrawIndirectParametersSize, kDrawIndexedIndirectParametersSize) + bufferOffset; + const indirectBuffer = t.device.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.INDIRECT, + }); + t.trackForCleanup(indirectBuffer); + + const outputTexture = t.device.createTexture({ + format: 'rgba8unorm', + size: [1, 1, 1], + usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.RENDER_ATTACHMENT, + }); + t.trackForCleanup(outputTexture); + + // Initialize outputTexture to green. + const encoder = t.device.createCommandEncoder(); + t.RecordInitializeTextureColor(encoder, outputTexture, { r: 0.0, g: 1.0, b: 0.0, a: 1.0 }); + + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: outputTexture.createView(), + loadOp: 'load', + storeOp: 'store', + }, + ], + }); + renderPass.setPipeline(renderPipeline); + + let indexBuffer = undefined; + if (test_indexed_draw) { + indexBuffer = t.device.createBuffer({ + size: 4, + usage: GPUBufferUsage.INDEX, + }); + t.trackForCleanup(indexBuffer); + renderPass.setIndexBuffer(indexBuffer, 'uint16'); + renderPass.drawIndexedIndirect(indirectBuffer, bufferOffset); + } else { + renderPass.drawIndirect(indirectBuffer, bufferOffset); + } + + renderPass.end(); + t.queue.submit([encoder.finish()]); + + // The indirect buffer should be lazily cleared to 0, so we actually draw nothing and the color + // attachment will keep its original color (green) after we end the render pass. + t.CheckBufferAndOutputTexture(indirectBuffer, bufferSize, outputTexture); + }); + +g.test('indirect_buffer_for_dispatch_indirect') + .desc( + `Verify when we use a GPUBuffer as an indirect buffer for dispatchWorkgroupsIndirect() just + after the creation of that GPUBuffer, all the contents in that GPUBuffer have been initialized + to 0.` + ) + .paramsSubcasesOnly(u => u.combine('bufferOffset', [0, 16])) + .fn(async t => { + const { bufferOffset } = t.params; + + const computePipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ + code: ` + @group(0) @binding(0) var outImage : texture_storage_2d<rgba8unorm, write>; + + @compute @workgroup_size(1) fn main() { + textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(1.0, 0.0, 0.0, 1.0)); + }`, + }), + entryPoint: 'main', + }, + }); + + const kDispatchIndirectParametersSize = 12; + const bufferSize = kDispatchIndirectParametersSize + bufferOffset; + const indirectBuffer = t.device.createBuffer({ + size: bufferSize, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.INDIRECT, + }); + t.trackForCleanup(indirectBuffer); + + const outputTexture = t.device.createTexture({ + format: 'rgba8unorm', + size: [1, 1, 1], + usage: + GPUTextureUsage.COPY_SRC | + GPUTextureUsage.RENDER_ATTACHMENT | + GPUTextureUsage.STORAGE_BINDING, + }); + t.trackForCleanup(outputTexture); + + // Initialize outputTexture to green. + const encoder = t.device.createCommandEncoder(); + t.RecordInitializeTextureColor(encoder, outputTexture, { r: 0.0, g: 1.0, b: 0.0, a: 1.0 }); + + const bindGroup = t.device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: outputTexture.createView(), + }, + ], + }); + + // The indirect buffer should be lazily cleared to 0, so we actually don't execute the compute + // shader and the output texture should keep its original color (green). + const computePass = encoder.beginComputePass(); + computePass.setBindGroup(0, bindGroup); + computePass.setPipeline(computePipeline); + computePass.dispatchWorkgroupsIndirect(indirectBuffer, bufferOffset); + computePass.end(); + t.queue.submit([encoder.finish()]); + + // The indirect buffer should be lazily cleared to 0, so we actually draw nothing and the color + // attachment will keep its original color (green) after we end the compute pass. + t.CheckBufferAndOutputTexture(indirectBuffer, bufferSize, outputTexture); + }); diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_copy.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_copy.ts new file mode 100644 index 0000000000..8f835e0f85 --- /dev/null +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_copy.ts @@ -0,0 +1,66 @@ +import { assert } from '../../../../../common/util/util.js'; +import { EncodableTextureFormat, kTextureFormatInfo } from '../../../../capability_info.js'; +import { virtualMipSize } from '../../../../util/texture/base.js'; +import { CheckContents } from '../texture_zero.spec.js'; + +export const checkContentsByBufferCopy: CheckContents = ( + t, + params, + texture, + state, + subresourceRange +) => { + for (const { level: mipLevel, layer } of subresourceRange.each()) { + assert(params.format in kTextureFormatInfo); + const format = params.format as EncodableTextureFormat; + + t.expectSingleColor(texture, format, { + size: [t.textureWidth, t.textureHeight, t.textureDepth], + dimension: params.dimension, + slice: layer, + layout: { mipLevel, aspect: params.aspect }, + exp: t.stateToTexelComponents[state], + }); + } +}; + +export const checkContentsByTextureCopy: CheckContents = ( + t, + params, + texture, + state, + subresourceRange +) => { + for (const { level, layer } of subresourceRange.each()) { + assert(params.format in kTextureFormatInfo); + const format = params.format as EncodableTextureFormat; + + const [width, height, depth] = virtualMipSize( + params.dimension, + [t.textureWidth, t.textureHeight, t.textureDepth], + level + ); + + const dst = t.device.createTexture({ + dimension: params.dimension, + size: [width, height, depth], + format: params.format, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.COPY_SRC, + }); + t.trackForCleanup(dst); + + const commandEncoder = t.device.createCommandEncoder(); + commandEncoder.copyTextureToTexture( + { texture, mipLevel: level, origin: { x: 0, y: 0, z: layer } }, + { texture: dst, mipLevel: 0 }, + { width, height, depthOrArrayLayers: depth } + ); + t.queue.submit([commandEncoder.finish()]); + + t.expectSingleColor(dst, format, { + size: [width, height, depth], + exp: t.stateToTexelComponents[state], + layout: { mipLevel: 0, aspect: params.aspect }, + }); + } +}; diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_ds_test.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_ds_test.ts new file mode 100644 index 0000000000..1851945e42 --- /dev/null +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_ds_test.ts @@ -0,0 +1,197 @@ +import { assert } from '../../../../../common/util/util.js'; +import { kTextureFormatInfo } from '../../../../capability_info.js'; +import { GPUTest } from '../../../../gpu_test.js'; +import { virtualMipSize } from '../../../../util/texture/base.js'; +import { CheckContents } from '../texture_zero.spec.js'; + +function makeFullscreenVertexModule(device: GPUDevice) { + return device.createShaderModule({ + code: ` + @vertex + fn main(@builtin(vertex_index) VertexIndex : u32) + -> @builtin(position) vec4<f32> { + var pos : array<vec2<f32>, 3> = array<vec2<f32>, 3>( + vec2<f32>(-1.0, -3.0), + vec2<f32>( 3.0, 1.0), + vec2<f32>(-1.0, 1.0)); + return vec4<f32>(pos[VertexIndex], 0.0, 1.0); + } + `, + }); +} + +function getDepthTestEqualPipeline( + t: GPUTest, + format: GPUTextureFormat, + sampleCount: number, + expected: number +): GPURenderPipeline { + return t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + entryPoint: 'main', + module: makeFullscreenVertexModule(t.device), + }, + fragment: { + entryPoint: 'main', + module: t.device.createShaderModule({ + code: ` + struct Outputs { + @builtin(frag_depth) FragDepth : f32, + @location(0) outSuccess : f32, + }; + + @fragment + fn main() -> Outputs { + var output : Outputs; + output.FragDepth = f32(${expected}); + output.outSuccess = 1.0; + return output; + } + `, + }), + targets: [{ format: 'r8unorm' }], + }, + depthStencil: { + format, + depthCompare: 'equal', + }, + primitive: { topology: 'triangle-list' }, + multisample: { count: sampleCount }, + }); +} + +function getStencilTestEqualPipeline( + t: GPUTest, + format: GPUTextureFormat, + sampleCount: number +): GPURenderPipeline { + return t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + entryPoint: 'main', + module: makeFullscreenVertexModule(t.device), + }, + fragment: { + entryPoint: 'main', + module: t.device.createShaderModule({ + code: ` + @fragment + fn main() -> @location(0) f32 { + return 1.0; + } + `, + }), + targets: [{ format: 'r8unorm' }], + }, + depthStencil: { + format, + stencilFront: { compare: 'equal' }, + stencilBack: { compare: 'equal' }, + }, + primitive: { topology: 'triangle-list' }, + multisample: { count: sampleCount }, + }); +} + +const checkContents: (type: 'depth' | 'stencil', ...args: Parameters<CheckContents>) => void = ( + type, + t, + params, + texture, + state, + subresourceRange +) => { + const formatInfo = kTextureFormatInfo[params.format]; + + assert(params.dimension === '2d'); + for (const viewDescriptor of t.generateTextureViewDescriptorsForRendering( + 'all', + subresourceRange + )) { + assert(viewDescriptor.baseMipLevel !== undefined); + const [width, height] = virtualMipSize( + params.dimension, + [t.textureWidth, t.textureHeight, 1], + viewDescriptor.baseMipLevel + ); + + const renderTexture = t.device.createTexture({ + size: [width, height, 1], + format: 'r8unorm', + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + sampleCount: params.sampleCount, + }); + + let resolveTexture = undefined; + let resolveTarget = undefined; + if (params.sampleCount > 1) { + resolveTexture = t.device.createTexture({ + size: [width, height, 1], + format: 'r8unorm', + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + }); + resolveTarget = resolveTexture.createView(); + } + + const commandEncoder = t.device.createCommandEncoder(); + commandEncoder.pushDebugGroup('checkContentsWithDepthStencil'); + + const pass = commandEncoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTexture.createView(), + resolveTarget, + clearValue: [0, 0, 0, 0], + loadOp: 'load', + storeOp: 'store', + }, + ], + depthStencilAttachment: { + view: texture.createView(viewDescriptor), + depthStoreOp: formatInfo.depth ? 'store' : undefined, + depthLoadOp: formatInfo.depth ? 'load' : undefined, + stencilStoreOp: formatInfo.stencil ? 'store' : undefined, + stencilLoadOp: formatInfo.stencil ? 'load' : undefined, + }, + }); + + switch (type) { + case 'depth': { + const expectedDepth = t.stateToTexelComponents[state].Depth; + assert(expectedDepth !== undefined); + + pass.setPipeline( + getDepthTestEqualPipeline(t, params.format, params.sampleCount, expectedDepth) + ); + break; + } + + case 'stencil': { + const expectedStencil = t.stateToTexelComponents[state].Stencil; + assert(expectedStencil !== undefined); + + pass.setPipeline(getStencilTestEqualPipeline(t, params.format, params.sampleCount)); + pass.setStencilReference(expectedStencil); + break; + } + } + + pass.draw(3); + pass.end(); + + commandEncoder.popDebugGroup(); + t.queue.submit([commandEncoder.finish()]); + + t.expectSingleColor(resolveTexture || renderTexture, 'r8unorm', { + size: [width, height, 1], + exp: { R: 1 }, + }); + } +}; + +export const checkContentsByDepthTest = (...args: Parameters<CheckContents>) => + checkContents('depth', ...args); + +export const checkContentsByStencilTest = (...args: Parameters<CheckContents>) => + checkContents('stencil', ...args); diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_sampling.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_sampling.ts new file mode 100644 index 0000000000..f739c128dc --- /dev/null +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/check_texture/by_sampling.ts @@ -0,0 +1,157 @@ +import { assert, unreachable } from '../../../../../common/util/util.js'; +import { EncodableTextureFormat, kTextureFormatInfo } from '../../../../capability_info.js'; +import { virtualMipSize } from '../../../../util/texture/base.js'; +import { + kTexelRepresentationInfo, + getSingleDataType, + getComponentReadbackTraits, +} from '../../../../util/texture/texel_data.js'; +import { CheckContents } from '../texture_zero.spec.js'; + +export const checkContentsBySampling: CheckContents = ( + t, + params, + texture, + state, + subresourceRange +) => { + assert(params.format in kTextureFormatInfo); + const format = params.format as EncodableTextureFormat; + const rep = kTexelRepresentationInfo[format]; + + for (const { level, layers } of subresourceRange.mipLevels()) { + const [width, height, depth] = virtualMipSize( + params.dimension, + [t.textureWidth, t.textureHeight, t.textureDepth], + level + ); + + const { ReadbackTypedArray, shaderType } = getComponentReadbackTraits( + getSingleDataType(format) + ); + + const componentOrder = rep.componentOrder; + const componentCount = componentOrder.length; + + // For single-component textures, generates .r + // For multi-component textures, generates ex.) + // .rgba[i], .bgra[i], .rgb[i] + const indexExpression = + componentCount === 1 + ? componentOrder[0].toLowerCase() + : componentOrder.map(c => c.toLowerCase()).join('') + '[i]'; + + const _xd = '_' + params.dimension; + const _multisampled = params.sampleCount > 1 ? '_multisampled' : ''; + const texelIndexExpression = + params.dimension === '2d' + ? 'vec2<i32>(GlobalInvocationID.xy)' + : params.dimension === '3d' + ? 'vec3<i32>(GlobalInvocationID.xyz)' + : params.dimension === '1d' + ? 'i32(GlobalInvocationID.x)' + : unreachable(); + const computePipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + entryPoint: 'main', + module: t.device.createShaderModule({ + code: ` + struct Constants { + level : i32 + }; + + @group(0) @binding(0) var<uniform> constants : Constants; + @group(0) @binding(1) var myTexture : texture${_multisampled}${_xd}<${shaderType}>; + + struct Result { + values : array<${shaderType}> + }; + @group(0) @binding(3) var<storage, read_write> result : Result; + + @compute @workgroup_size(1) + fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) { + let flatIndex : u32 = ${componentCount}u * ( + ${width}u * ${height}u * GlobalInvocationID.z + + ${width}u * GlobalInvocationID.y + + GlobalInvocationID.x + ); + let texel : vec4<${shaderType}> = textureLoad( + myTexture, ${texelIndexExpression}, constants.level); + + for (var i : u32 = 0u; i < ${componentCount}u; i = i + 1u) { + result.values[flatIndex + i] = texel.${indexExpression}; + } + }`, + }), + }, + }); + + for (const layer of layers) { + const ubo = t.device.createBuffer({ + mappedAtCreation: true, + size: 4, + usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + }); + new Int32Array(ubo.getMappedRange(), 0, 1)[0] = level; + ubo.unmap(); + + const byteLength = + width * height * depth * ReadbackTypedArray.BYTES_PER_ELEMENT * rep.componentOrder.length; + const resultBuffer = t.device.createBuffer({ + size: byteLength, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + t.trackForCleanup(resultBuffer); + + const bindGroup = t.device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { buffer: ubo }, + }, + { + binding: 1, + resource: texture.createView({ + baseArrayLayer: layer, + arrayLayerCount: 1, + dimension: params.dimension, + }), + }, + { + binding: 3, + resource: { + buffer: resultBuffer, + }, + }, + ], + }); + + const commandEncoder = t.device.createCommandEncoder(); + const pass = commandEncoder.beginComputePass(); + pass.setPipeline(computePipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(width, height, depth); + pass.end(); + t.queue.submit([commandEncoder.finish()]); + ubo.destroy(); + + const expectedValues = new ReadbackTypedArray(new ArrayBuffer(byteLength)); + const expectedState = t.stateToTexelComponents[state]; + let i = 0; + for (let d = 0; d < depth; ++d) { + for (let h = 0; h < height; ++h) { + for (let w = 0; w < width; ++w) { + for (const c of rep.componentOrder) { + const value = expectedState[c]; + assert(value !== undefined); + expectedValues[i++] = value; + } + } + } + } + t.expectGPUBufferValuesEqual(resultBuffer, expectedValues); + } + } +}; diff --git a/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/texture_zero.spec.ts b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/texture_zero.spec.ts new file mode 100644 index 0000000000..cdb383ad65 --- /dev/null +++ b/dom/webgpu/tests/cts/checkout/src/webgpu/api/operation/resource_init/texture_zero.spec.ts @@ -0,0 +1,645 @@ +export const description = ` +Test uninitialized textures are initialized to zero when read. + +TODO: +- test by sampling depth/stencil [1] +- test by copying out of stencil [2] +- test compressed texture formats [3] +`; + +// MAINTENANCE_TODO: This is a test file, it probably shouldn't export anything. +// Everything that's exported should be moved to another file. + +import { TestCaseRecorder, TestParams } from '../../../../common/framework/fixture.js'; +import { + kUnitCaseParamsBuilder, + ParamTypeOf, +} from '../../../../common/framework/params_builder.js'; +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { assert, unreachable } from '../../../../common/util/util.js'; +import { + kTextureFormatInfo, + kTextureAspects, + kUncompressedTextureFormats, + EncodableTextureFormat, + UncompressedTextureFormat, + textureDimensionAndFormatCompatible, + kTextureDimensions, +} from '../../../capability_info.js'; +import { GPUConst } from '../../../constants.js'; +import { GPUTest, GPUTestSubcaseBatchState } from '../../../gpu_test.js'; +import { virtualMipSize } from '../../../util/texture/base.js'; +import { createTextureUploadBuffer } from '../../../util/texture/layout.js'; +import { BeginEndRange, SubresourceRange } from '../../../util/texture/subresource.js'; +import { PerTexelComponent, kTexelRepresentationInfo } from '../../../util/texture/texel_data.js'; + +export enum UninitializeMethod { + Creation = 'Creation', // The texture was just created. It is uninitialized. + StoreOpClear = 'StoreOpClear', // The texture was rendered to with GPUStoreOp "clear" +} +const kUninitializeMethods = Object.keys(UninitializeMethod) as UninitializeMethod[]; + +export const enum ReadMethod { + Sample = 'Sample', // The texture is sampled from + CopyToBuffer = 'CopyToBuffer', // The texture is copied to a buffer + CopyToTexture = 'CopyToTexture', // The texture is copied to another texture + DepthTest = 'DepthTest', // The texture is read as a depth buffer + StencilTest = 'StencilTest', // The texture is read as a stencil buffer + ColorBlending = 'ColorBlending', // Read the texture by blending as a color attachment + Storage = 'Storage', // Read the texture as a storage texture +} + +// Test with these mip level counts +type MipLevels = 1 | 5; +const kMipLevelCounts: MipLevels[] = [1, 5]; + +// For each mip level count, define the mip ranges to leave uninitialized. +const kUninitializedMipRangesToTest: { [k in MipLevels]: BeginEndRange[] } = { + 1: [{ begin: 0, end: 1 }], // Test the only mip + 5: [ + { begin: 0, end: 2 }, + { begin: 3, end: 4 }, + ], // Test a range and a single mip +}; + +// Test with these sample counts. +const kSampleCounts: number[] = [1, 4]; + +// Test with these layer counts. +type LayerCounts = 1 | 7; + +// For each layer count, define the layers to leave uninitialized. +const kUninitializedLayerRangesToTest: { [k in LayerCounts]: BeginEndRange[] } = { + 1: [{ begin: 0, end: 1 }], // Test the only layer + 7: [ + { begin: 2, end: 4 }, + { begin: 6, end: 7 }, + ], // Test a range and a single layer +}; + +// Enums to abstract over color / depth / stencil values in textures. Depending on the texture format, +// the data for each value may have a different representation. These enums are converted to a +// representation such that their values can be compared. ex.) An integer is needed to upload to an +// unsigned normalized format, but its value is read as a float in the shader. +export const enum InitializedState { + Canary, // Set on initialized subresources. It should stay the same. On discarded resources, we should observe zero. + Zero, // We check that uninitialized subresources are in this state when read back. +} + +const initializedStateAsFloat = { + [InitializedState.Zero]: 0, + [InitializedState.Canary]: 1, +}; + +const initializedStateAsUint = { + [InitializedState.Zero]: 0, + [InitializedState.Canary]: 1, +}; + +const initializedStateAsSint = { + [InitializedState.Zero]: 0, + [InitializedState.Canary]: -1, +}; + +function initializedStateAsColor( + state: InitializedState, + format: GPUTextureFormat +): [number, number, number, number] { + let value; + if (format.indexOf('uint') !== -1) { + value = initializedStateAsUint[state]; + } else if (format.indexOf('sint') !== -1) { + value = initializedStateAsSint[state]; + } else { + value = initializedStateAsFloat[state]; + } + return [value, value, value, value]; +} + +const initializedStateAsDepth = { + [InitializedState.Zero]: 0, + [InitializedState.Canary]: 0.8, +}; + +const initializedStateAsStencil = { + [InitializedState.Zero]: 0, + [InitializedState.Canary]: 42, +}; + +function getRequiredTextureUsage( + format: UncompressedTextureFormat, + sampleCount: number, + uninitializeMethod: UninitializeMethod, + readMethod: ReadMethod +): GPUTextureUsageFlags { + let usage: GPUTextureUsageFlags = GPUConst.TextureUsage.COPY_DST; + + switch (uninitializeMethod) { + case UninitializeMethod.Creation: + break; + case UninitializeMethod.StoreOpClear: + usage |= GPUConst.TextureUsage.RENDER_ATTACHMENT; + break; + default: + unreachable(); + } + + switch (readMethod) { + case ReadMethod.CopyToBuffer: + case ReadMethod.CopyToTexture: + usage |= GPUConst.TextureUsage.COPY_SRC; + break; + case ReadMethod.Sample: + usage |= GPUConst.TextureUsage.TEXTURE_BINDING; + break; + case ReadMethod.Storage: + usage |= GPUConst.TextureUsage.STORAGE_BINDING; + break; + case ReadMethod.DepthTest: + case ReadMethod.StencilTest: + case ReadMethod.ColorBlending: + usage |= GPUConst.TextureUsage.RENDER_ATTACHMENT; + break; + default: + unreachable(); + } + + if (sampleCount > 1) { + // Copies to multisampled textures are not allowed. We need OutputAttachment to initialize + // canary data in multisampled textures. + usage |= GPUConst.TextureUsage.RENDER_ATTACHMENT; + } + + if (!kTextureFormatInfo[format].copyDst) { + // Copies are not possible. We need OutputAttachment to initialize + // canary data. + assert(kTextureFormatInfo[format].renderable); + usage |= GPUConst.TextureUsage.RENDER_ATTACHMENT; + } + + return usage; +} + +export class TextureZeroInitTest extends GPUTest { + readonly stateToTexelComponents: { [k in InitializedState]: PerTexelComponent<number> }; + + private p: TextureZeroParams; + constructor(sharedState: GPUTestSubcaseBatchState, rec: TestCaseRecorder, params: TestParams) { + super(sharedState, rec, params); + this.p = params as TextureZeroParams; + + const stateToTexelComponents = (state: InitializedState) => { + const [R, G, B, A] = initializedStateAsColor(state, this.p.format); + return { + R, + G, + B, + A, + Depth: initializedStateAsDepth[state], + Stencil: initializedStateAsStencil[state], + }; + }; + + this.stateToTexelComponents = { + [InitializedState.Zero]: stateToTexelComponents(InitializedState.Zero), + [InitializedState.Canary]: stateToTexelComponents(InitializedState.Canary), + }; + } + + get textureWidth(): number { + let width = 1 << this.p.mipLevelCount; + if (this.p.nonPowerOfTwo) { + width = 2 * width - 1; + } + return width; + } + + get textureHeight(): number { + if (this.p.dimension === '1d') { + return 1; + } + + let height = 1 << this.p.mipLevelCount; + if (this.p.nonPowerOfTwo) { + height = 2 * height - 1; + } + return height; + } + + get textureDepth(): number { + return this.p.dimension === '3d' ? 11 : 1; + } + + get textureDepthOrArrayLayers(): number { + return this.p.dimension === '2d' ? this.p.layerCount : this.textureDepth; + } + + // Used to iterate subresources and check that their uninitialized contents are zero when accessed + *iterateUninitializedSubresources(): Generator<SubresourceRange> { + for (const mipRange of kUninitializedMipRangesToTest[this.p.mipLevelCount]) { + for (const layerRange of kUninitializedLayerRangesToTest[this.p.layerCount]) { + yield new SubresourceRange({ mipRange, layerRange }); + } + } + } + + // Used to iterate and initialize other subresources not checked for zero-initialization. + // Zero-initialization of uninitialized subresources should not have side effects on already + // initialized subresources. + *iterateInitializedSubresources(): Generator<SubresourceRange> { + const uninitialized: boolean[][] = new Array(this.p.mipLevelCount); + for (let level = 0; level < uninitialized.length; ++level) { + uninitialized[level] = new Array(this.p.layerCount); + } + for (const subresources of this.iterateUninitializedSubresources()) { + for (const { level, layer } of subresources.each()) { + uninitialized[level][layer] = true; + } + } + for (let level = 0; level < uninitialized.length; ++level) { + for (let layer = 0; layer < uninitialized[level].length; ++layer) { + if (!uninitialized[level][layer]) { + yield new SubresourceRange({ + mipRange: { begin: level, count: 1 }, + layerRange: { begin: layer, count: 1 }, + }); + } + } + } + } + + *generateTextureViewDescriptorsForRendering( + aspect: GPUTextureAspect, + subresourceRange?: SubresourceRange + ): Generator<GPUTextureViewDescriptor> { + const viewDescriptor: GPUTextureViewDescriptor = { + dimension: '2d', + aspect, + }; + + if (subresourceRange === undefined) { + return viewDescriptor; + } + + for (const { level, layer } of subresourceRange.each()) { + yield { + ...viewDescriptor, + baseMipLevel: level, + mipLevelCount: 1, + baseArrayLayer: layer, + arrayLayerCount: 1, + }; + } + } + + private initializeWithStoreOp( + state: InitializedState, + texture: GPUTexture, + subresourceRange?: SubresourceRange + ): void { + const commandEncoder = this.device.createCommandEncoder(); + commandEncoder.pushDebugGroup('initializeWithStoreOp'); + + for (const viewDescriptor of this.generateTextureViewDescriptorsForRendering( + 'all', + subresourceRange + )) { + if (kTextureFormatInfo[this.p.format].color) { + commandEncoder + .beginRenderPass({ + colorAttachments: [ + { + view: texture.createView(viewDescriptor), + storeOp: 'store', + clearValue: initializedStateAsColor(state, this.p.format), + loadOp: 'clear', + }, + ], + }) + .end(); + } else { + const depthStencilAttachment: GPURenderPassDepthStencilAttachment = { + view: texture.createView(viewDescriptor), + }; + if (kTextureFormatInfo[this.p.format].depth) { + depthStencilAttachment.depthClearValue = initializedStateAsDepth[state]; + depthStencilAttachment.depthLoadOp = 'clear'; + depthStencilAttachment.depthStoreOp = 'store'; + } + if (kTextureFormatInfo[this.p.format].stencil) { + depthStencilAttachment.stencilClearValue = initializedStateAsStencil[state]; + depthStencilAttachment.stencilLoadOp = 'clear'; + depthStencilAttachment.stencilStoreOp = 'store'; + } + commandEncoder + .beginRenderPass({ + colorAttachments: [], + depthStencilAttachment, + }) + .end(); + } + } + + commandEncoder.popDebugGroup(); + this.queue.submit([commandEncoder.finish()]); + } + + private initializeWithCopy( + texture: GPUTexture, + state: InitializedState, + subresourceRange: SubresourceRange + ): void { + assert(this.p.format in kTextureFormatInfo); + const format = this.p.format as EncodableTextureFormat; + + const firstSubresource = subresourceRange.each().next().value; + assert(typeof firstSubresource !== 'undefined'); + + const [largestWidth, largestHeight, largestDepth] = virtualMipSize( + this.p.dimension, + [this.textureWidth, this.textureHeight, this.textureDepth], + firstSubresource.level + ); + + const rep = kTexelRepresentationInfo[format]; + const texelData = new Uint8Array(rep.pack(rep.encode(this.stateToTexelComponents[state]))); + const { buffer, bytesPerRow, rowsPerImage } = createTextureUploadBuffer( + texelData, + this.device, + format, + this.p.dimension, + [largestWidth, largestHeight, largestDepth] + ); + + const commandEncoder = this.device.createCommandEncoder(); + + for (const { level, layer } of subresourceRange.each()) { + const [width, height, depth] = virtualMipSize( + this.p.dimension, + [this.textureWidth, this.textureHeight, this.textureDepth], + level + ); + + commandEncoder.copyBufferToTexture( + { + buffer, + bytesPerRow, + rowsPerImage, + }, + { texture, mipLevel: level, origin: { x: 0, y: 0, z: layer } }, + { width, height, depthOrArrayLayers: depth } + ); + } + this.queue.submit([commandEncoder.finish()]); + buffer.destroy(); + } + + initializeTexture( + texture: GPUTexture, + state: InitializedState, + subresourceRange: SubresourceRange + ): void { + if (this.p.sampleCount > 1 || !kTextureFormatInfo[this.p.format].copyDst) { + // Copies to multisampled textures not yet specified. + // Use a storeOp for now. + assert(kTextureFormatInfo[this.p.format].renderable); + this.initializeWithStoreOp(state, texture, subresourceRange); + } else { + this.initializeWithCopy(texture, state, subresourceRange); + } + } + + discardTexture(texture: GPUTexture, subresourceRange: SubresourceRange): void { + const commandEncoder = this.device.createCommandEncoder(); + commandEncoder.pushDebugGroup('discardTexture'); + + for (const desc of this.generateTextureViewDescriptorsForRendering('all', subresourceRange)) { + if (kTextureFormatInfo[this.p.format].color) { + commandEncoder + .beginRenderPass({ + colorAttachments: [ + { + view: texture.createView(desc), + storeOp: 'discard', + loadOp: 'load', + }, + ], + }) + .end(); + } else { + const depthStencilAttachment: GPURenderPassDepthStencilAttachment = { + view: texture.createView(desc), + }; + if (kTextureFormatInfo[this.p.format].depth) { + depthStencilAttachment.depthLoadOp = 'load'; + depthStencilAttachment.depthStoreOp = 'discard'; + } + if (kTextureFormatInfo[this.p.format].stencil) { + depthStencilAttachment.stencilLoadOp = 'load'; + depthStencilAttachment.stencilStoreOp = 'discard'; + } + commandEncoder + .beginRenderPass({ + colorAttachments: [], + depthStencilAttachment, + }) + .end(); + } + } + + commandEncoder.popDebugGroup(); + this.queue.submit([commandEncoder.finish()]); + } +} + +const kTestParams = kUnitCaseParamsBuilder + .combine('dimension', kTextureDimensions) + .combine('readMethod', [ + ReadMethod.CopyToBuffer, + ReadMethod.CopyToTexture, + ReadMethod.Sample, + ReadMethod.DepthTest, + ReadMethod.StencilTest, + ]) + // [3] compressed formats + .combine('format', kUncompressedTextureFormats) + .filter(({ dimension, format }) => textureDimensionAndFormatCompatible(dimension, format)) + .beginSubcases() + .combine('aspect', kTextureAspects) + .unless(({ readMethod, format, aspect }) => { + const info = kTextureFormatInfo[format]; + return ( + (readMethod === ReadMethod.DepthTest && (!info.depth || aspect === 'stencil-only')) || + (readMethod === ReadMethod.StencilTest && (!info.stencil || aspect === 'depth-only')) || + (readMethod === ReadMethod.ColorBlending && !info.color) || + // [1]: Test with depth/stencil sampling + (readMethod === ReadMethod.Sample && (info.depth || info.stencil)) || + (aspect === 'depth-only' && !info.depth) || + (aspect === 'stencil-only' && !info.stencil) || + (aspect === 'all' && info.depth && info.stencil) || + // Cannot copy from a packed depth format. + // [2]: Test copying out of the stencil aspect. + ((readMethod === ReadMethod.CopyToBuffer || readMethod === ReadMethod.CopyToTexture) && + (format === 'depth24plus' || format === 'depth24plus-stencil8')) + ); + }) + .combine('mipLevelCount', kMipLevelCounts) + // 1D texture can only have a single mip level + .unless(p => p.dimension === '1d' && p.mipLevelCount !== 1) + .combine('sampleCount', kSampleCounts) + .unless( + ({ readMethod, sampleCount }) => + // We can only read from multisampled textures by sampling. + sampleCount > 1 && + (readMethod === ReadMethod.CopyToBuffer || readMethod === ReadMethod.CopyToTexture) + ) + // Multisampled textures may only have one mip + .unless(({ sampleCount, mipLevelCount }) => sampleCount > 1 && mipLevelCount > 1) + .combine('uninitializeMethod', kUninitializeMethods) + .unless(({ dimension, readMethod, uninitializeMethod, format, sampleCount }) => { + const formatInfo = kTextureFormatInfo[format]; + return ( + dimension !== '2d' && + (sampleCount > 1 || + formatInfo.depth || + formatInfo.stencil || + readMethod === ReadMethod.DepthTest || + readMethod === ReadMethod.StencilTest || + readMethod === ReadMethod.ColorBlending || + uninitializeMethod === UninitializeMethod.StoreOpClear) + ); + }) + .expandWithParams(function* ({ dimension }) { + switch (dimension) { + case '2d': + yield { layerCount: 1 as LayerCounts }; + yield { layerCount: 7 as LayerCounts }; + break; + case '1d': + case '3d': + yield { layerCount: 1 as LayerCounts }; + break; + } + }) + // Multisampled 3D / 2D array textures not supported. + .unless(({ sampleCount, layerCount }) => sampleCount > 1 && layerCount > 1) + .unless(({ format, sampleCount, uninitializeMethod, readMethod }) => { + const usage = getRequiredTextureUsage(format, sampleCount, uninitializeMethod, readMethod); + const info = kTextureFormatInfo[format]; + + return ( + ((usage & GPUConst.TextureUsage.RENDER_ATTACHMENT) !== 0 && !info.renderable) || + ((usage & GPUConst.TextureUsage.STORAGE_BINDING) !== 0 && !info.storage) || + (sampleCount > 1 && !info.multisample) + ); + }) + .combine('nonPowerOfTwo', [false, true]) + .combine('canaryOnCreation', [false, true]) + .filter(({ canaryOnCreation, format }) => { + // We can only initialize the texture if it's encodable or renderable. + const canInitialize = format in kTextureFormatInfo || kTextureFormatInfo[format].renderable; + + // Filter out cases where we want canary values but can't initialize. + return !canaryOnCreation || canInitialize; + }); + +type TextureZeroParams = ParamTypeOf<typeof kTestParams>; + +export type CheckContents = ( + t: TextureZeroInitTest, + params: TextureZeroParams, + texture: GPUTexture, + state: InitializedState, + subresourceRange: SubresourceRange +) => void; + +import { checkContentsByBufferCopy, checkContentsByTextureCopy } from './check_texture/by_copy.js'; +import { + checkContentsByDepthTest, + checkContentsByStencilTest, +} from './check_texture/by_ds_test.js'; +import { checkContentsBySampling } from './check_texture/by_sampling.js'; + +const checkContentsImpl: { [k in ReadMethod]: CheckContents } = { + Sample: checkContentsBySampling, + CopyToBuffer: checkContentsByBufferCopy, + CopyToTexture: checkContentsByTextureCopy, + DepthTest: checkContentsByDepthTest, + StencilTest: checkContentsByStencilTest, + ColorBlending: t => t.skip('Not implemented'), + Storage: t => t.skip('Not implemented'), +}; + +export const g = makeTestGroup(TextureZeroInitTest); + +g.test('uninitialized_texture_is_zero') + .params(kTestParams) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase(kTextureFormatInfo[t.params.format].feature); + }) + .fn(async t => { + const usage = getRequiredTextureUsage( + t.params.format, + t.params.sampleCount, + t.params.uninitializeMethod, + t.params.readMethod + ); + + const texture = t.device.createTexture({ + size: [t.textureWidth, t.textureHeight, t.textureDepthOrArrayLayers], + format: t.params.format, + dimension: t.params.dimension, + usage, + mipLevelCount: t.params.mipLevelCount, + sampleCount: t.params.sampleCount, + }); + t.trackForCleanup(texture); + + if (t.params.canaryOnCreation) { + // Initialize some subresources with canary values + for (const subresourceRange of t.iterateInitializedSubresources()) { + t.initializeTexture(texture, InitializedState.Canary, subresourceRange); + } + } + + switch (t.params.uninitializeMethod) { + case UninitializeMethod.Creation: + break; + case UninitializeMethod.StoreOpClear: + // Initialize the rest of the resources. + for (const subresourceRange of t.iterateUninitializedSubresources()) { + t.initializeTexture(texture, InitializedState.Canary, subresourceRange); + } + // Then use a store op to discard their contents. + for (const subresourceRange of t.iterateUninitializedSubresources()) { + t.discardTexture(texture, subresourceRange); + } + break; + default: + unreachable(); + } + + // Check that all uninitialized resources are zero. + for (const subresourceRange of t.iterateUninitializedSubresources()) { + checkContentsImpl[t.params.readMethod]( + t, + t.params, + texture, + InitializedState.Zero, + subresourceRange + ); + } + + if (t.params.canaryOnCreation) { + // Check the all other resources are unchanged. + for (const subresourceRange of t.iterateInitializedSubresources()) { + checkContentsImpl[t.params.readMethod]( + t, + t.params, + texture, + InitializedState.Canary, + subresourceRange + ); + } + } + }); |