From f3c4e19700456b0f55d1737190252c2e35ff8bc1 Mon Sep 17 00:00:00 2001 From: Greggman Date: Tue, 6 Aug 2024 10:23:41 -0700 Subject: [PATCH] Add WGSL textureLoad tests for storage formats (#3890) --- src/webgpu/listing_meta.json | 5 +- .../call/builtin/textureLoad.spec.ts | 569 ++++++++---------- 2 files changed, 258 insertions(+), 316 deletions(-) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 2d100ac6b6a7..2616e8f976db 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1561,7 +1561,10 @@ "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_1d:*": { "subcaseMS": 83.312 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_2d:*": { "subcaseMS": 96.737 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:sampled_3d:*": { "subcaseMS": 158.534 }, - "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_texel_formats:*": { "subcaseMS": 471.569 }, + "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_1d:*": { "subcaseMS": 41.569 }, + "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_2d:*": { "subcaseMS": 41.569 }, + "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_2d_array:*": { "subcaseMS": 41.569 }, + "webgpu:shader,execution,expression,call,builtin,textureLoad:storage_textures_3d:*": { "subcaseMS": 41.569 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:arrayed:*": { "subcaseMS": 8.102 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:sampled:*": { "subcaseMS": 2.101 }, "webgpu:shader,execution,expression,call,builtin,textureNumLayers:storage:*": { "subcaseMS": 8.000 }, diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts index 1a7818c76379..4103feef6ac7 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts @@ -22,7 +22,6 @@ TODO: Test stencil8 format. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { unreachable, iterRange } from '../../../../../../common/util/util.js'; import { canUseAsRenderTarget, isCompressedFloatTextureFormat, @@ -37,13 +36,6 @@ import { textureDimensionAndFormatCompatible, } from '../../../../../format_info.js'; import { GPUTest } from '../../../../../gpu_test.js'; -import { - kFloat32Format, - kFloat16Format, - numberToFloatBits, - pack4x8unorm, - pack4x8snorm, -} from '../../../../../util/conversion.js'; import { maxMipLevelCount, virtualMipSize } from '../../../../../util/texture/base.js'; import { TexelFormats } from '../../../../types.js'; @@ -574,105 +566,26 @@ Parameters: t.expectOK(res); }); -// Returns texel values to use as inputs for textureLoad. -// Values are kept simple to avoid rounding issues. -function shaderValues(format: string, type: string) { - switch (type) { - case 'f32': { - switch (format) { - case 'rbga8snorm': - // prettier-ignore - return [ - { r: 0.0, g: 0.0, b: 0.0, a: 0.0, }, - { r: 0.2, g: 0.4, b: 0.6, a: 0.8, }, - { r: -0.2, g: -0.4, b: -0.6, a: -0.8, }, - { r: 0.2, g: -0.4, b: 0.6, a: -0.8, }, - { r: -0.2, g: 0.4, b: -0.6, a: 0.8, }, - { r: 0.2, g: 0.2, b: 0.2, a: 0.2, }, - { r: -0.2, g: -0.2, b: -0.2, a: -0.2, }, - { r: 0.4, g: 0.4, b: 0.4, a: 0.4, }, - { r: -0.4, g: -0.4, b: -0.4, a: -0.4, }, - { r: 0.6, g: 0.6, b: 0.6, a: 0.6, }, - { r: -0.6, g: -0.6, b: -0.6, a: -0.6, }, - { r: 0.8, g: 0.8, b: 0.8, a: 0.8, }, - { r: -0.8, g: -0.8, b: -0.8, a: -0.8, }, - ]; - case 'rgba8unorm': - case 'bgra8unorm': - // prettier-ignore - return [ - { r: 0.0, g: 0.0, b: 0.0, a: 0.0, }, - { r: 0.2, g: 0.4, b: 0.6, a: 0.8, }, - { r: 0.9, g: 0.4, b: 0.6, a: 0.8, }, - { r: 0.2, g: 0.9, b: 0.6, a: 0.8, }, - { r: 0.2, g: 0.4, b: 0.9, a: 0.8, }, - { r: 0.2, g: 0.4, b: 0.6, a: 0.9, }, - { r: 0.2, g: 0.2, b: 0.2, a: 0.2, }, - { r: 0.4, g: 0.4, b: 0.4, a: 0.4, }, - { r: 0.6, g: 0.6, b: 0.6, a: 0.6, }, - { r: 0.8, g: 0.8, b: 0.8, a: 0.8, }, - ]; - default: - // Stick within 16-bit ranges. - // prettier-ignore - return [ - { r: 100, g: 128, b: 100, a: 128, }, - { r: 64, g: 32, b: 32, a: 64, }, - { r: 8, g: 0, b: 8, a: 0, }, - { r: 0, g: 0, b: 0, a: 0, }, - { r: -100, g: 128, b: 100, a: 128, }, - { r: -64, g: 32, b: 32, a: 64, }, - { r: -8, g: 0, b: 8, a: 0, }, - { r: 100, g: -128, b: 100, a: 128, }, - { r: 64, g: -32, b: 32, a: 64, }, - { r: 8, g: 0, b: 8, a: 0, }, - { r: 100, g: 128, b: -100, a: 128, }, - { r: 64, g: 32, b: -32, a: 64, }, - { r: 8, g: 0, b: -8, a: 0, }, - { r: 100, g: 128, b: 100, a: -128, }, - { r: 64, g: 32, b: 32, a: -64, }, - { r: 8, g: 0, b: 8, a: 0, }, - ]; - } - break; - } - case 'u32': - // Keep all ranges within u8. - // prettier-ignore - return [ - { r: 0, g: 0, b: 0, a: 0, }, - { r: 0, g: 8, b: 16, a: 128, }, - { r: 8, g: 16, b: 32, a: 64, }, - { r: 16, g: 32, b: 64, a: 128, }, - { r: 255, g: 254, b: 253, a: 252, }, - { r: 255, g: 255, b: 255, a: 255, }, - { r: 128, g: 64, b: 32, a: 16, }, - { r: 64, g: 32, b: 16, a: 8, }, - { r: 32, g: 16, b: 8, a: 0, }, - ]; - case 'i32': - // Keep all ranges i8 - // prettier-ignore - return [ - { r: 0, g: 0, b: 0, a: 0, }, - { r: 0, g: -8, b: 16, a: 127, }, - { r: 8, g: 16, b: -32, a: 64, }, - { r: -16, g: 32, b: 64, a: -128, }, - { r: 127, g: 126, b: 125, a: 124, }, - { r: -128, g: -127, b: -126, a: -125, }, - { r: 127, g: 127, b: 127, a: 127, }, - { r: -128, g: -128, b: -128, a: -128, }, - ]; - default: - unreachable(`unhandled shader type ${type}`); - break; - } - return []; -} +g.test('storage_textures_1d') + .specURL('https://www.w3.org/TR/WGSL/#textureload') + .desc( + ` +C is i32 or u32 -g.test('storage_texel_formats') - .desc('Test loading of texel formats') - .params(u => u.combineWithParams([...TexelFormats, { format: 'bgra8unorm', _shaderType: 'f32' }])) +fn textureLoad(t: texture_storage_1d, coords: C) -> vec4 + +Parameters: + * t: The sampled texture to read from + * coords: The 0-based texel coordinate +` + ) + .params(u => + u + .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('C', ['i32', 'u32'] as const) + ) .beforeAllSubcases(t => { t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures')); if (t.params.format === 'bgra8unorm') { @@ -681,223 +594,249 @@ g.test('storage_texel_formats') t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); } }) - .fn(t => { - const { format, _shaderType } = t.params; - const values = shaderValues(format, _shaderType); - - // To avoid rounding issues, unorm and snorm values are repacked in the shader. - let useType = _shaderType; - let assignValue = `v`; - if (format === 'bgra8unorm' || format === 'rgba8unorm') { - useType = 'u32'; - assignValue = `vec4u(pack4x8unorm(v),0,0,0)`; - } else if (format === 'rgba8snorm') { - useType = 'u32'; - assignValue = `vec4u(pack4x8snorm(v),0,0,0)`; - } - const wgsl = ` -requires readonly_and_readwrite_storage_textures; - -@group(0) @binding(0) -var tex : texture_storage_1d<${format}, read>; - -@group(0) @binding(1) -var out : array>; - -@compute @workgroup_size(${values.length}) -fn main(@builtin(global_invocation_id) gid : vec3u) { - let v = textureLoad(tex, gid.x); - out[gid.x] = ${assignValue}; -}`; - - const bytesPerRow = 256; - let bytesPerTexel = 4; - switch (format) { - case 'rgba16uint': - case 'rgba16sint': - case 'rgba16float': - case 'rg32uint': - case 'rg32sint': - case 'rg32float': - bytesPerTexel = 8; - break; - case 'rgba32uint': - case 'rgba32sint': - case 'rgba32float': - bytesPerTexel = 16; - break; - default: - break; - } + .fn(async t => { + const { format, samplePoints, C } = t.params; - const textureSize: GPUExtent3D = { - width: bytesPerRow / bytesPerTexel, - height: 1, - depthOrArrayLayers: 1, - }; - const texture = t.createTextureTracked({ - format: format as GPUTextureFormat, + // We want at least 3 blocks or something wide enough for 3 mip levels. + const [width] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const size = [width, 1]; + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, dimension: '1d', - size: textureSize, - mipLevelCount: 1, - usage: GPUTextureUsage.STORAGE_BINDING | GPUTextureUsage.COPY_DST, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + + const calls: TextureCall[] = generateTextureBuiltinInputs1D(50, { + method: samplePoints, + descriptor, + hashInputs: [format, samplePoints, C], + }).map(({ coords }) => { + return { + builtin: 'textureLoad', + coordType: C === 'i32' ? 'i' : 'u', + coords: normalizedCoordToTexelLoadTestCoord(descriptor, 0, C, coords), + }; }); - const outputBuffer = t.makeBufferWithContents( - new Uint32Array([...iterRange(values.length * 4, x => 0)]), - GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE + const textureType = `texture_storage_1d<${format}, read>`; + const viewDescriptor = {}; + const sampler = undefined; + const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const res = await checkCallResults( + t, + { texels, descriptor, viewDescriptor }, + textureType, + sampler, + calls, + results ); - t.trackForCleanup(outputBuffer); - - const transformed = values.flatMap(x => { - switch (format) { - case 'rgba8unorm': - return pack4x8unorm(x.r, x.g, x.b, x.a); - case 'bgra8unorm': - return pack4x8unorm(x.b, x.g, x.r, x.a); - case 'rgba8snorm': - return pack4x8snorm(x.r, x.g, x.b, x.a); - case 'r32uint': - case 'r32sint': - return x.r; - case 'rg32uint': - case 'rg32sint': - return [x.r, x.g]; - case 'rgba32uint': - case 'rgba32sint': - return [x.r, x.g, x.b, x.a]; - case 'rgba8uint': - case 'rgba8sint': - return (x.r & 0xff) | ((x.g & 0xff) << 8) | ((x.b & 0xff) << 16) | ((x.a & 0xff) << 24); - case 'rgba16uint': - case 'rgba16sint': - return [(x.r & 0xffff) | ((x.g & 0xffff) << 16), (x.b & 0xffff) | ((x.a & 0xffff) << 16)]; - case 'r32float': - return numberToFloatBits(x.r, kFloat32Format); - case 'rg32float': - return [numberToFloatBits(x.r, kFloat32Format), numberToFloatBits(x.g, kFloat32Format)]; - case 'rgba32float': - return [ - numberToFloatBits(x.r, kFloat32Format), - numberToFloatBits(x.g, kFloat32Format), - numberToFloatBits(x.b, kFloat32Format), - numberToFloatBits(x.a, kFloat32Format), - ]; - case 'rgba16float': - return [ - (numberToFloatBits(x.r, kFloat16Format) & 0xffff) | - ((numberToFloatBits(x.g, kFloat16Format) & 0xffff) << 16), - (numberToFloatBits(x.b, kFloat16Format) & 0xffff) | - ((numberToFloatBits(x.a, kFloat16Format) & 0xffff) << 16), - ]; - default: - unreachable(`unhandled format ${format}`); - break; - } - return 0; - }); + t.expectOK(res); + }); - const texelBuffer = t.makeBufferWithContents( - new Uint32Array([ - ...iterRange(bytesPerRow, x => { - if (x < transformed.length) { - return transformed[x]; - } else { - return 0; - } - }), - ]), - GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE - ); - t.trackForCleanup(texelBuffer); - - const pipeline = t.device.createComputePipeline({ - layout: 'auto', - compute: { - module: t.device.createShaderModule({ - code: wgsl, - }), - entryPoint: 'main', - }, - }); - const bg = t.device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { - binding: 0, - resource: texture.createView({ - format: format as GPUTextureFormat, - dimension: '1d', - }), - }, - { - binding: 1, - resource: { - buffer: outputBuffer, - }, - }, - ], +g.test('storage_textures_2d') + .specURL('https://www.w3.org/TR/WGSL/#textureload') + .desc( + ` +C is i32 or u32 + +fn textureLoad(t: texture_storage_2d, coords: vec2) -> vec4 + +Parameters: + * t: The sampled texture to read from + * coords: The 0-based texel coordinate +` + ) + .params(u => + u + .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('C', ['i32', 'u32'] as const) + ) + .beforeAllSubcases(t => { + t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures')); + if (t.params.format === 'bgra8unorm') { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + } else { + t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); + } + }) + .fn(async t => { + const { format, samplePoints, C } = t.params; + + // We want at least 3 blocks or something wide enough for 3 mip levels. + const size = chooseTextureSize({ minSize: 8, minBlocks: 3, format }); + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + descriptor, + hashInputs: [format, samplePoints, C], + }).map(({ coords }) => { + return { + builtin: 'textureLoad', + coordType: C === 'i32' ? 'i' : 'u', + coords: normalizedCoordToTexelLoadTestCoord(descriptor, 0, C, coords), + }; }); + const textureType = `texture_storage_2d<${format}, read>`; + const viewDescriptor = {}; + const sampler = undefined; + const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const res = await checkCallResults( + t, + { texels, descriptor, viewDescriptor }, + textureType, + sampler, + calls, + results + ); + t.expectOK(res); + }); + +g.test('storage_textures_2d_array') + .specURL('https://www.w3.org/TR/WGSL/#textureload') + .desc( + ` +C is i32 or u32 +A is i32 or u32 + +fn textureLoad(t: texture_storage_2d, coords: vec2, array_index: A) -> vec4 + +Parameters: + * t: The sampled texture to read from + * coords: The 0-based texel coordinate + * array_index: The 0-based texture array index +` + ) + .params(u => + u + .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('C', ['i32', 'u32'] as const) + .combine('A', ['i32', 'u32'] as const) + ) + .beforeAllSubcases(t => { + t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures')); + if (t.params.format === 'bgra8unorm') { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + } else { + t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); + } + }) + .fn(async t => { + const { format, samplePoints, C, A } = t.params; + + // We want at least 3 blocks or something wide enough for 3 mip levels. + const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); - const encoder = t.device.createCommandEncoder(); - encoder.copyBufferToTexture( - { - buffer: texelBuffer, - offset: 0, - bytesPerRow, - rowsPerImage: 1, - }, - { texture }, - textureSize + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + descriptor, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + hashInputs: [format, samplePoints, C, A], + }).map(({ coords, arrayIndex }) => { + return { + builtin: 'textureLoad', + coordType: C === 'i32' ? 'i' : 'u', + coords: normalizedCoordToTexelLoadTestCoord(descriptor, 0, C, coords), + arrayIndexType: A === 'i32' ? 'i' : 'u', + arrayIndex, + }; + }); + const textureType = `texture_storage_2d_array<${format}, read>`; + const viewDescriptor: GPUTextureViewDescriptor = { + dimension: '2d-array', + }; + const sampler = undefined; + const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const res = await checkCallResults( + t, + { texels, descriptor, viewDescriptor }, + textureType, + sampler, + calls, + results ); + t.expectOK(res); + }); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bg); - pass.dispatchWorkgroups(1, 1, 1); - pass.end(); - t.queue.submit([encoder.finish()]); - - const expected = new Uint32Array( - values.flatMap(x => { - switch (format) { - case 'r32uint': - case 'r32sint': - return [x.r, 0, 0, 1]; - case 'rg32uint': - case 'rg32sint': - return [x.r, x.g, 0, 1]; - case 'r32float': - return [ - numberToFloatBits(x.r, kFloat32Format), - 0, - 0, - numberToFloatBits(1, kFloat32Format), - ]; - case 'rg32float': - return [ - numberToFloatBits(x.r, kFloat32Format), - numberToFloatBits(x.g, kFloat32Format), - 0, - numberToFloatBits(1, kFloat32Format), - ]; - case 'rgba32float': - case 'rgba16float': - return [ - numberToFloatBits(x.r, kFloat32Format), - numberToFloatBits(x.g, kFloat32Format), - numberToFloatBits(x.b, kFloat32Format), - numberToFloatBits(x.a, kFloat32Format), - ]; - case 'rgba8unorm': - case 'bgra8unorm': - return [pack4x8unorm(x.r, x.g, x.b, x.a), 0, 0, 0]; - case 'rgba8snorm': - return [pack4x8snorm(x.r, x.g, x.b, x.a), 0, 0, 0]; - default: - break; - } - return [x.r, x.g, x.b, x.a]; - }) +g.test('storage_textures_3d') + .specURL('https://www.w3.org/TR/WGSL/#textureload') + .desc( + ` +C is i32 or u32 + +fn textureLoad(t: texture_storage_2d, coords: vec3) -> vec4 + +Parameters: + * t: The sampled texture to read from + * coords: The 0-based texel coordinate +` + ) + .params(u => + u + .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('C', ['i32', 'u32'] as const) + ) + .beforeAllSubcases(t => { + t.skipIf(!t.hasLanguageFeature('readonly_and_readwrite_storage_textures')); + if (t.params.format === 'bgra8unorm') { + t.selectDeviceOrSkipTestCase('bgra8unorm-storage'); + } else { + t.skipIfTextureFormatNotUsableAsStorageTexture(t.params.format as GPUTextureFormat); + } + }) + .fn(async t => { + const { format, samplePoints, C } = t.params; + + // We want at least 3 blocks or something wide enough for 3 mip levels. + const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.STORAGE_BINDING, + dimension: '3d', + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + + const calls: TextureCall[] = generateTextureBuiltinInputs3D(50, { + method: samplePoints, + descriptor, + hashInputs: [format, samplePoints, C], + }).map(({ coords }) => { + return { + builtin: 'textureLoad', + coordType: C === 'i32' ? 'i' : 'u', + coords: normalizedCoordToTexelLoadTestCoord(descriptor, 0, C, coords), + }; + }); + const textureType = `texture_storage_3d<${format}, read>`; + const viewDescriptor = {}; + const sampler = undefined; + const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const res = await checkCallResults( + t, + { texels, descriptor, viewDescriptor }, + textureType, + sampler, + calls, + results ); - t.expectGPUBufferValuesEqual(outputBuffer, expected); + t.expectOK(res); });