From 40867c3c83526e442350dfca7bbb8ee332a3431e Mon Sep 17 00:00:00 2001 From: Greggman Date: Wed, 8 Jan 2025 04:41:11 +0900 Subject: [PATCH] Compat: refactor storage_texture tests (#4099) 1. Skip them if there are no storage textures 2. Refactor them so they don't need storage buffers to test. Also did some minor JavaScript styling. I hope that's okay. While the CTS doesn't have a style guide, JS is generally methods, properties, and variable names are camelCase and classes are CamelCase. --- .../storage_texture/read_only.spec.ts | 276 +++++++++--------- .../storage_texture/read_write.spec.ts | 34 ++- 2 files changed, 157 insertions(+), 153 deletions(-) diff --git a/src/webgpu/api/operation/storage_texture/read_only.spec.ts b/src/webgpu/api/operation/storage_texture/read_only.spec.ts index b5224eec2b9d..769036751258 100644 --- a/src/webgpu/api/operation/storage_texture/read_only.spec.ts +++ b/src/webgpu/api/operation/storage_texture/read_only.spec.ts @@ -15,7 +15,7 @@ import { kColorTextureFormats, kTextureFormatInfo, } from '../../../format_info.js'; -import { GPUTest } from '../../../gpu_test.js'; +import { GPUTest, MaxLimitsTestMixin } from '../../../gpu_test.js'; import { kValidShaderStages, TValidShaderStage } from '../../../util/shader.js'; function ComponentCount(format: ColorTextureFormat): number { @@ -47,7 +47,7 @@ function ComponentCount(format: ColorTextureFormat): number { } class F extends GPUTest { - InitTextureAndGetExpectedOutputBufferData( + initTextureAndGetExpectedOutputBufferData( storageTexture: GPUTexture, format: ColorTextureFormat ): ArrayBuffer { @@ -59,10 +59,10 @@ class F extends GPUTest { const depthOrArrayLayers = storageTexture.depthOrArrayLayers; const texelData = new ArrayBuffer(bytesPerBlock * width * height * depthOrArrayLayers); - const texelTypedDataView = this.GetTypedArrayBufferViewForTexelData(texelData, format); + const texelTypedDataView = this.getTypedArrayBufferViewForTexelData(texelData, format); const componentCount = ComponentCount(format); const outputBufferData = new ArrayBuffer(4 * 4 * width * height * depthOrArrayLayers); - const outputBufferTypedData = this.GetTypedArrayBufferForOutputBufferData( + const outputBufferTypedData = this.getTypedArrayBufferForOutputBufferData( outputBufferData, format ); @@ -174,7 +174,7 @@ class F extends GPUTest { return outputBufferData; } - GetTypedArrayBufferForOutputBufferData(arrayBuffer: ArrayBuffer, format: ColorTextureFormat) { + getTypedArrayBufferForOutputBufferData(arrayBuffer: ArrayBuffer, format: ColorTextureFormat) { switch (kTextureFormatInfo[format].color.type) { case 'uint': return new Uint32Array(arrayBuffer); @@ -186,7 +186,7 @@ class F extends GPUTest { } } - GetTypedArrayBufferViewForTexelData(arrayBuffer: ArrayBuffer, format: ColorTextureFormat) { + getTypedArrayBufferViewForTexelData(arrayBuffer: ArrayBuffer, format: ColorTextureFormat) { switch (format) { case 'r32uint': case 'rg32uint': @@ -219,7 +219,7 @@ class F extends GPUTest { } } - GetOutputBufferWGSLType(format: ColorTextureFormat) { + getOutputBufferWGSLType(format: ColorTextureFormat) { switch (kTextureFormatInfo[format].color.type) { case 'uint': return 'vec4u'; @@ -234,7 +234,7 @@ class F extends GPUTest { } } - DoTransform( + doTransform( storageTexture: GPUTexture, shaderStage: TValidShaderStage, format: ColorTextureFormat, @@ -256,23 +256,22 @@ class F extends GPUTest { const textureDeclaration = ` @group(0) @binding(0) var readOnlyTexture: ${declaration}<${format}, read>; `; - const bindingResourceDeclaration = ` - ${textureDeclaration} - @group(0) @binding(1) - var outputBuffer : array<${this.GetOutputBufferWGSLType(format)}>; - `; const bindGroupEntries = [ { binding: 0, resource: storageTexture.createView(), }, - { - binding: 1, - resource: { - buffer: outputBuffer, - }, - }, + ...(shaderStage === 'compute' + ? [ + { + binding: 1, + resource: { + buffer: outputBuffer, + }, + }, + ] + : []), ]; const commandEncoder = this.device.createCommandEncoder(); @@ -296,7 +295,10 @@ class F extends GPUTest { } const computeShader = ` - ${bindingResourceDeclaration} + ${textureDeclaration} + @group(0) @binding(1) + var outputBuffer : array<${this.getOutputBufferWGSLType(format)}>; + @compute @workgroup_size( ${storageTexture.width}, ${storageTexture.height}, ${storageTexture.depthOrArrayLayers}) @@ -334,61 +336,42 @@ class F extends GPUTest { break; case '2d': textureLoadCoord = - storageTexture.depthOrArrayLayers > 1 ? 'textureCoord, z' : 'textureCoord'; + storageTexture.depthOrArrayLayers > 1 ? 'textureCoord, coordZ' : 'textureCoord'; break; case '3d': - textureLoadCoord = 'vec3u(textureCoord, z)'; + textureLoadCoord = 'vec3u(textureCoord, coordZ)'; break; } - const fragmentShader = ` - ${bindingResourceDeclaration} + const shader = ` + ${textureDeclaration} @fragment - fn main(@builtin(position) fragCoord: vec4f) -> @location(0) vec4f { - let textureCoord = vec2u(fragCoord.xy); - let storageTextureTexelCountPerImage = ${storageTexture.width * storageTexture.height}u; - for (var z = 0u; z < ${storageTexture.depthOrArrayLayers}; z++) { - let initialValue = textureLoad(readOnlyTexture, ${textureLoadCoord}); - let outputIndex = - storageTextureTexelCountPerImage * z + textureCoord.y * ${storageTexture.width} + - textureCoord.x; - outputBuffer[outputIndex] = initialValue; - } - return vec4f(0.0, 1.0, 0.0, 1.0); - }`; - const vertexShader = ` - @vertex - fn main(@builtin(vertex_index) vertexIndex : u32) -> @builtin(position) vec4f { - var pos = array( - vec2f(-1.0, -1.0), - vec2f(-1.0, 1.0), - vec2f( 1.0, -1.0), - vec2f(-1.0, 1.0), - vec2f( 1.0, -1.0), - vec2f( 1.0, 1.0)); - return vec4f(pos[vertexIndex], 0.0, 1.0); - } - `; + fn fs(@builtin(position) fragCoord: vec4f) -> @location(0) vec4u { + let coordX = u32(fragCoord.x); + let coordY = u32(fragCoord.y) % ${storageTexture.height}u; + let coordZ = u32(fragCoord.y) / ${storageTexture.height}u; + let textureCoord = vec2u(coordX, coordY); + return bitcast(textureLoad(readOnlyTexture, ${textureLoadCoord})); + } + + @vertex + fn vs(@builtin(vertex_index) vertexIndex : u32) -> @builtin(position) vec4f { + var pos = array( + vec2f(-1.0, 3.0), + vec2f( 3.0, -1.0), + vec2f(-1.0, -1.0)); + return vec4f(pos[vertexIndex], 0.0, 1.0); + } + `; + + const module = this.device.createShaderModule({ + code: shader, + }); const renderPipeline = this.device.createRenderPipeline({ layout: 'auto', - vertex: { - module: this.device.createShaderModule({ - code: vertexShader, - }), - }, - fragment: { - module: this.device.createShaderModule({ - code: fragmentShader, - }), - targets: [ - { - format: 'rgba8unorm', - }, - ], - }, - primitive: { - topology: 'triangle-list', - }, + vertex: { module }, + fragment: { module, targets: [{ format: 'rgba32uint' }] }, + primitive: { topology: 'triangle-list' }, }); const bindGroup = this.device.createBindGroup({ @@ -396,10 +379,14 @@ class F extends GPUTest { entries: bindGroupEntries, }); + // This is just so our buffer compare is the same as the compute stage. + // Otherwise, we'd have to pad every row to a multiple of 256 bytes and + // change the comparison code to take that into account. + assert(storageTexture.width === 16, `width must be 16 because we require 256 bytesPerRow`); const placeholderColorTexture = this.createTextureTracked({ - size: [storageTexture.width, storageTexture.height, 1], - usage: GPUTextureUsage.RENDER_ATTACHMENT, - format: 'rgba8unorm', + size: [storageTexture.width, storageTexture.height * storageTexture.depthOrArrayLayers], + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + format: 'rgba32uint', }); const renderPassEncoder = commandEncoder.beginRenderPass({ @@ -407,54 +394,50 @@ class F extends GPUTest { { view: placeholderColorTexture.createView(), loadOp: 'clear', - clearValue: { r: 0, g: 0, b: 0, a: 0 }, storeOp: 'store', }, ], }); renderPassEncoder.setPipeline(renderPipeline); renderPassEncoder.setBindGroup(0, bindGroup); - renderPassEncoder.draw(6); + renderPassEncoder.draw(3); renderPassEncoder.end(); + + commandEncoder.copyTextureToBuffer( + { texture: placeholderColorTexture }, + { + buffer: outputBuffer, + bytesPerRow: 256, + }, + placeholderColorTexture + ); break; } case 'vertex': { - // For each texel location (coordX, coordY), draw one point at (coordX + 0.5, coordY + 0.5) - // in the storageTexture.width * storageTexture.height grid, and save all the texel values - // at (coordX, coordY, z) (z >= 0 && z < storageTexture.depthOrArrayLayers) into the - // corresponding vertex shader outputs. - let vertexOutputs = ''; - for (let layer = 0; layer < storageTexture.depthOrArrayLayers; ++layer) { - vertexOutputs = vertexOutputs.concat( - ` - @location(${layer + 1}) @interpolate(flat, either) - vertex_out${layer}: ${this.GetOutputBufferWGSLType(format)},` - ); - } - + // We draw storageTexture.Width by (storageTexture.height * storageTexture.depthOrArrayLayers) + // points via 'point-list' to a placeholderColorTexture of the same size. + // + // We use the @builtin(vertex_index) to compute a coord in the source texture + // and use that same coord to compute a place to render in the point in the placeholder. let loadFromTextureWGSL = ''; switch (storageTexture.dimension) { case '1d': loadFromTextureWGSL = ` - output.vertex_out0 = textureLoad(readOnlyTexture, coordX);`; + output.vertex_out = textureLoad(readOnlyTexture, coordX);`; break; case '2d': if (storageTexture.depthOrArrayLayers === 1) { loadFromTextureWGSL = ` - output.vertex_out0 = textureLoad(readOnlyTexture, vec2u(coordX, coordY));`; + output.vertex_out = textureLoad(readOnlyTexture, vec2u(coordX, coordY));`; } else { - for (let z = 0; z < storageTexture.depthOrArrayLayers; ++z) { - loadFromTextureWGSL = loadFromTextureWGSL.concat(` - output.vertex_out${z} = - textureLoad(readOnlyTexture, vec2u(coordX, coordY), ${z});`); - } + loadFromTextureWGSL = loadFromTextureWGSL.concat(` + output.vertex_out = + textureLoad(readOnlyTexture, vec2u(coordX, coordY), coordZ);`); } break; case '3d': - for (let z = 0; z < storageTexture.depthOrArrayLayers; ++z) { - loadFromTextureWGSL = loadFromTextureWGSL.concat(` - output.vertex_out${z} = textureLoad(readOnlyTexture, vec3u(coordX, coordY, ${z}));`); - } + loadFromTextureWGSL = loadFromTextureWGSL.concat(` + output.vertex_out = textureLoad(readOnlyTexture, vec3u(coordX, coordY, coordZ));`); break; } @@ -470,57 +453,39 @@ class F extends GPUTest { } const shader = ` - ${bindingResourceDeclaration} + ${textureDeclaration} struct VertexOutput { @builtin(position) my_pos: vec4f, - @location(0) @interpolate(flat, either) tex_coord: vec2u, - ${vertexOutputs} + @location(0) @interpolate(flat, either) + vertex_out: ${this.getOutputBufferWGSLType(format)}, } @vertex fn vs_main(@builtin(vertex_index) vertexIndex : u32) -> VertexOutput { var output : VertexOutput; let coordX = vertexIndex % ${storageTexture.width}u; - let coordY = vertexIndex / ${storageTexture.width}u; - // Each vertex in the mesh take an even step along X axis from -1.0 to 1.0. - let posXStep = f32(${2.0 / storageTexture.width}); - // As well as along Y axis. - let posYStep = f32(${2.0 / storageTexture.height}); - // And the vertex located in the middle of the step, i.e. with a bias of 0.5 step. - let outputPosX = -1.0 + posXStep * 0.5 + posXStep * f32(coordX); - let outputPosY = -1.0 + posYStep * 0.5 + posYStep * f32(coordY); - output.my_pos = vec4f(outputPosX, outputPosY, 0.0, 1.0); - output.tex_coord = vec2u(coordX, coordY); + let coordY = vertexIndex / ${storageTexture.width}u % ${storageTexture.height}u; + let coordZ = vertexIndex / ${storageTexture.width * storageTexture.height}u; + let writePos = vec2f(f32(coordX), f32(coordY + coordZ * ${storageTexture.height})); + let destSize = vec2f( + ${storageTexture.width}, + ${storageTexture.height * storageTexture.depthOrArrayLayers}); + output.my_pos = vec4f((((writePos + 0.5) / destSize) * 2.0 - 1.0) * vec2f(1, -1), 0.0, 1.0); ${loadFromTextureWGSL} return output; } @fragment - fn fs_main(fragmentInput : VertexOutput) -> @location(0) vec4f { - let storageTextureTexelCountPerImage = ${storageTexture.width * storageTexture.height}u; - ${outputToBufferWGSL} - return vec4f(0.0, 1.0, 0.0, 1.0); + fn fs_main(fragmentInput : VertexOutput) -> @location(0) vec4u { + let v = fragmentInput.vertex_out; + return bitcast(v); } `; + const module = this.device.createShaderModule({ code: shader }); const renderPipeline = this.device.createRenderPipeline({ layout: 'auto', - vertex: { - module: this.device.createShaderModule({ - code: shader, - }), - }, - fragment: { - module: this.device.createShaderModule({ - code: shader, - }), - targets: [ - { - format: 'rgba8unorm', - }, - ], - }, - primitive: { - topology: 'point-list', - }, + vertex: { module }, + fragment: { module, targets: [{ format: 'rgba32uint' }] }, + primitive: { topology: 'point-list' }, }); const bindGroup = this.device.createBindGroup({ @@ -529,9 +494,9 @@ class F extends GPUTest { }); const placeholderColorTexture = this.createTextureTracked({ - size: [storageTexture.width, storageTexture.height, 1], - usage: GPUTextureUsage.RENDER_ATTACHMENT, - format: 'rgba8unorm', + size: [storageTexture.width, storageTexture.height * storageTexture.depthOrArrayLayers], + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + format: 'rgba32uint', }); const renderPassEncoder = commandEncoder.beginRenderPass({ @@ -546,8 +511,19 @@ class F extends GPUTest { }); renderPassEncoder.setPipeline(renderPipeline); renderPassEncoder.setBindGroup(0, bindGroup); - renderPassEncoder.draw(storageTexture.width * storageTexture.height); + const texelCount = + storageTexture.width * storageTexture.height * storageTexture.depthOrArrayLayers; + renderPassEncoder.draw(texelCount); renderPassEncoder.end(); + + commandEncoder.copyTextureToBuffer( + { texture: placeholderColorTexture }, + { + buffer: outputBuffer, + bytesPerRow: 256, + }, + placeholderColorTexture + ); break; } } @@ -556,7 +532,7 @@ class F extends GPUTest { } } -export const g = makeTestGroup(F); +export const g = makeTestGroup(MaxLimitsTestMixin(F)); g.test('basic') .desc( @@ -586,7 +562,23 @@ g.test('basic') .fn(t => { const { format, shaderStage, dimension, depthOrArrayLayers } = t.params; - const kWidth = 8; + if (t.isCompatibility) { + if (shaderStage === 'fragment') { + t.skipIf( + !(t.device.limits.maxStorageTexturesInFragmentStage! > 0), + `maxStorageTexturesInFragmentStage(${t.device.limits + .maxStorageTexturesInFragmentStage!}) is not > 0` + ); + } else if (shaderStage === 'vertex') { + t.skipIf( + !(t.device.limits.maxStorageTexturesInVertexStage! > 0), + `maxStorageTexturesInVertexStage(${t.device.limits + .maxStorageTexturesInVertexStage!}) is not > 0` + ); + } + } + + const kWidth = 16; const height = dimension === '1d' ? 1 : 8; const storageTexture = t.createTextureTracked({ format, @@ -595,14 +587,16 @@ g.test('basic') usage: GPUTextureUsage.COPY_SRC | GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, }); - const expectedData = t.InitTextureAndGetExpectedOutputBufferData(storageTexture, format); + const expectedData = t.initTextureAndGetExpectedOutputBufferData(storageTexture, format); + const bytesPerRow = 4 * 4 * kWidth; + assert(bytesPerRow === 256, 'bytesPerRow === 256'); const outputBuffer = t.createBufferTracked({ - size: 4 * 4 * kWidth * height * depthOrArrayLayers, - usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE, + size: bytesPerRow * height * depthOrArrayLayers, + usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, }); - t.DoTransform(storageTexture, shaderStage, format, outputBuffer); + t.doTransform(storageTexture, shaderStage, format, outputBuffer); switch (kTextureFormatInfo[format].color.type) { case 'uint': diff --git a/src/webgpu/api/operation/storage_texture/read_write.spec.ts b/src/webgpu/api/operation/storage_texture/read_write.spec.ts index 03f613284238..ad70d68352c4 100644 --- a/src/webgpu/api/operation/storage_texture/read_write.spec.ts +++ b/src/webgpu/api/operation/storage_texture/read_write.spec.ts @@ -9,7 +9,7 @@ import { makeTestGroup } from '../../../../common/framework/test_group.js'; import { assert, unreachable } from '../../../../common/util/util.js'; import { kTextureDimensions } from '../../../capability_info.js'; import { kColorTextureFormats, kTextureFormatInfo } from '../../../format_info.js'; -import { GPUTest } from '../../../gpu_test.js'; +import { GPUTest, MaxLimitsTestMixin } from '../../../gpu_test.js'; import { align } from '../../../util/math.js'; const kShaderStagesForReadWriteStorageTexture = ['fragment', 'compute'] as const; @@ -17,7 +17,7 @@ type ShaderStageForReadWriteStorageTexture = (typeof kShaderStagesForReadWriteStorageTexture)[number]; class F extends GPUTest { - GetInitialData(storageTexture: GPUTexture): ArrayBuffer { + getInitialData(storageTexture: GPUTexture): ArrayBuffer { const format = storageTexture.format; const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock; assert(bytesPerBlock !== undefined); @@ -26,7 +26,7 @@ class F extends GPUTest { const height = storageTexture.height; const depthOrArrayLayers = storageTexture.depthOrArrayLayers; const initialData = new ArrayBuffer(bytesPerBlock * width * height * depthOrArrayLayers); - const initialTypedData = this.GetTypedArrayBuffer(initialData, format); + const initialTypedData = this.getTypedArrayBuffer(initialData, format); for (let z = 0; z < depthOrArrayLayers; ++z) { for (let y = 0; y < height; ++y) { for (let x = 0; x < width; ++x) { @@ -48,7 +48,7 @@ class F extends GPUTest { return initialData; } - GetTypedArrayBuffer(arrayBuffer: ArrayBuffer, format: GPUTextureFormat) { + getTypedArrayBuffer(arrayBuffer: ArrayBuffer, format: GPUTextureFormat) { switch (format) { case 'r32sint': return new Int32Array(arrayBuffer); @@ -62,7 +62,7 @@ class F extends GPUTest { } } - GetExpectedData( + getExpectedData( shaderStage: ShaderStageForReadWriteStorageTexture, storageTexture: GPUTexture, initialData: ArrayBuffer @@ -80,8 +80,8 @@ class F extends GPUTest { const expectedData = new ArrayBuffer( bytesPerRowAlignment * (height * depthOrArrayLayers - 1) + bytesPerBlock * width ); - const expectedTypedData = this.GetTypedArrayBuffer(expectedData, format); - const initialTypedData = this.GetTypedArrayBuffer(initialData, format); + const expectedTypedData = this.getTypedArrayBuffer(expectedData, format); + const initialTypedData = this.getTypedArrayBuffer(initialData, format); for (let z = 0; z < depthOrArrayLayers; ++z) { for (let y = 0; y < height; ++y) { for (let x = 0; x < width; ++x) { @@ -110,7 +110,7 @@ class F extends GPUTest { return expectedData; } - RecordCommandsToTransform( + recordCommandsToTransform( device: GPUDevice, shaderStage: ShaderStageForReadWriteStorageTexture, commandEncoder: GPUCommandEncoder, @@ -298,7 +298,7 @@ class F extends GPUTest { } } -export const g = makeTestGroup(F); +export const g = makeTestGroup(MaxLimitsTestMixin(F)); g.test('basic') .desc( @@ -321,6 +321,16 @@ g.test('basic') .fn(t => { const { format, shaderStage, textureDimension, depthOrArrayLayers } = t.params; + if (t.isCompatibility) { + if (shaderStage === 'fragment') { + t.skipIf( + !(t.device.limits.maxStorageTexturesInFragmentStage! > 0), + `maxStorageTexturesInFragmentStage(${t.device.limits + .maxStorageTexturesInFragmentStage!}) is not > 0` + ); + } + } + // In compatibility mode the lowest maxComputeInvocationsPerWorkgroup is 128 vs non-compat which is 256 // So in non-compat we get 16 * 8 * 2, vs compat where we get 8 * 8 * 2 const kWidth = t.isCompatibility ? 8 : 16; @@ -334,7 +344,7 @@ g.test('basic') }); const bytesPerBlock = kTextureFormatInfo[format].bytesPerBlock; - const initialData = t.GetInitialData(storageTexture); + const initialData = t.getInitialData(storageTexture); t.queue.writeTexture( { texture: storageTexture }, initialData, @@ -347,9 +357,9 @@ g.test('basic') const commandEncoder = t.device.createCommandEncoder(); - t.RecordCommandsToTransform(t.device, shaderStage, commandEncoder, storageTexture); + t.recordCommandsToTransform(t.device, shaderStage, commandEncoder, storageTexture); - const expectedData = t.GetExpectedData(shaderStage, storageTexture, initialData); + const expectedData = t.getExpectedData(shaderStage, storageTexture, initialData); const readbackBuffer = t.createBufferTracked({ size: expectedData.byteLength, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,