From 955975e3476084a6e342fea6a80438ac1e4db293 Mon Sep 17 00:00:00 2001 From: Gregg Tavares Date: Wed, 30 Oct 2024 10:57:59 -0700 Subject: [PATCH] Fix readTextureToTexelViews to handle all types of formats and viewDimensions Issue #4017 --- src/webgpu/format_info.ts | 11 +- .../call/builtin/texture_utils.spec.ts | 109 ++++++++++++ .../expression/call/builtin/texture_utils.ts | 164 +++++++++++++----- 3 files changed, 242 insertions(+), 42 deletions(-) create mode 100644 src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts diff --git a/src/webgpu/format_info.ts b/src/webgpu/format_info.ts index 792a58f546b..5d3b5becf56 100644 --- a/src/webgpu/format_info.ts +++ b/src/webgpu/format_info.ts @@ -1822,12 +1822,21 @@ export function isRegularTextureFormat(format: GPUTextureFormat) { } /** - * Returns true of format is both compressed and a float format, for example 'bc6h-rgb-ufloat'. + * Returns true if format is both compressed and a float format, for example 'bc6h-rgb-ufloat'. */ export function isCompressedFloatTextureFormat(format: GPUTextureFormat) { return isCompressedTextureFormat(format) && format.includes('float'); } +/** + * Returns true if format is sint or uint + */ +export function isSintOrUintFormat(format: GPUTextureFormat) { + const info = kTextureFormatInfo[format]; + const type = info.color?.type ?? info.depth?.type ?? info.stencil?.type; + return type === 'sint' || type === 'uint'; +} + /** * Returns true of format can be multisampled. */ diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts new file mode 100644 index 00000000000..f4be99ee65b --- /dev/null +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.spec.ts @@ -0,0 +1,109 @@ +export const description = ` +Tests for texture_utils.ts +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { assert } from '../../../../../../common/util/util.js'; +import { GPUTest } from '../../../../../gpu_test.js'; +import { getTextureDimensionFromView, virtualMipSize } from '../../../../../util/texture/base.js'; +import { + kTexelRepresentationInfo, + PerTexelComponent, + TexelRepresentationInfo, +} from '../../../../../util/texture/texel_data.js'; + +import { + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + isSupportedViewFormatCombo, + readTextureToTexelViews, + texelsApproximatelyEqual, +} from './texture_utils.js'; + +export const g = makeTestGroup(GPUTest); + +function texelFormat(texel: Readonly>, rep: TexelRepresentationInfo) { + return rep.componentOrder.map(component => `${component}: ${texel[component]}`).join(', '); +} + +g.test('readTextureToTexelViews') + .desc('test readTextureToTexelViews for various formats and dimensions') + .params(u => + u + .combineWithParams([ + { srcFormat: 'r8unorm', texelViewFormat: 'rgba32float' }, + { srcFormat: 'r8sint', texelViewFormat: 'rgba32sint' }, + { srcFormat: 'r8uint', texelViewFormat: 'rgba32uint' }, + { srcFormat: 'rgba32float', texelViewFormat: 'rgba32float' }, + { srcFormat: 'rgba32uint', texelViewFormat: 'rgba32uint' }, + { srcFormat: 'rgba32sint', texelViewFormat: 'rgba32sint' }, + { srcFormat: 'depth24plus', texelViewFormat: 'rgba32float' }, + { srcFormat: 'depth24plus', texelViewFormat: 'r32float' }, + { srcFormat: 'depth24plus-stencil8', texelViewFormat: 'r32float' }, + { srcFormat: 'stencil8', texelViewFormat: 'rgba32sint' }, + ] as const) + .combine('viewDimension', ['1d', '2d', '2d-array', '3d', 'cube', 'cube-array'] as const) + .filter(t => isSupportedViewFormatCombo(t.srcFormat, t.viewDimension)) + ) + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); + }) + .fn(async t => { + const { srcFormat, texelViewFormat, viewDimension } = t.params; + const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format: srcFormat, viewDimension }); + const descriptor: GPUTextureDescriptor = { + format: srcFormat, + dimension: getTextureDimensionFromView(viewDimension), + size, + mipLevelCount: viewDimension === '1d' ? 1 : 3, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + }; + const { texels: expectedTexelViews, texture } = await createTextureWithRandomDataAndGetTexels( + t, + descriptor + ); + const actualTexelViews = await readTextureToTexelViews(t, texture, descriptor, texelViewFormat); + + assert(actualTexelViews.length === expectedTexelViews.length, 'num mip levels match'); + + const errors = []; + for (let mipLevel = 0; mipLevel < actualTexelViews.length; ++mipLevel) { + const actualMipLevelTexelView = actualTexelViews[mipLevel]; + const expectedMipLevelTexelView = expectedTexelViews[mipLevel]; + const mipLevelSize = virtualMipSize(texture.dimension, size, mipLevel); + + const actualRep = kTexelRepresentationInfo[actualMipLevelTexelView.format]; + const expectedRep = kTexelRepresentationInfo[expectedMipLevelTexelView.format]; + + for (let z = 0; z < mipLevelSize[2]; ++z) { + for (let y = 0; y < mipLevelSize[1]; ++y) { + for (let x = 0; x < mipLevelSize[0]; ++x) { + const actual = actualMipLevelTexelView.color({ x, y, z }); + const expected = expectedMipLevelTexelView.color({ x, y, z }); + // This currently expects the exact same values in actual vs expected. + // It's possible this needs to be relaxed slightly but only for non-integer formats. + // For now, if the tests pass everywhere, we'll keep it at 0 tolerance. + const maxFractionalDiff = 0; + if ( + !texelsApproximatelyEqual( + actual, + actualMipLevelTexelView.format, + expected, + expectedMipLevelTexelView.format, + maxFractionalDiff + ) + ) { + const actualStr = texelFormat(actual, actualRep); + const expectedStr = texelFormat(expected, expectedRep); + errors.push( + `texel at ${x}, ${y}, ${z}, expected: ${expectedStr}, actual: ${actualStr}` + ); + } + } + } + } + + assert(errors.length === 0, errors.join('\n')); + } + }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts index 72268c5b368..57c422cb675 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -8,6 +8,7 @@ import { isDepthOrStencilTextureFormat, isDepthTextureFormat, isEncodableTextureFormat, + isSintOrUintFormat, isStencilTextureFormat, kEncodableTextureFormats, kTextureFormatInfo, @@ -79,8 +80,8 @@ export function isSupportedViewFormatCombo( viewDimension: GPUTextureViewDimension ) { return !( - (isCompressedTextureFormat(format) || isDepthTextureFormat(format)) && - viewDimension === '3d' + (isCompressedTextureFormat(format) || isDepthOrStencilTextureFormat(format)) && + (viewDimension === '3d' || viewDimension === '1d') ); } @@ -928,6 +929,11 @@ function getMinAndMaxTexelValueForComponent( * or something similar to TexelView. */ export function getTexelViewFormatForTextureFormat(format: GPUTextureFormat) { + if (format.endsWith('sint')) { + return 'rgba32sint'; + } else if (format.endsWith('uint')) { + return 'rgba32uint'; + } return format.endsWith('-srgb') ? 'rgba8unorm-srgb' : 'rgba32float'; } @@ -1251,6 +1257,9 @@ const builtinNeedsDerivatives = (builtin: TextureBuiltin) => const isCubeViewDimension = (viewDescriptor?: GPUTextureViewDescriptor) => viewDescriptor?.dimension === 'cube' || viewDescriptor?.dimension === 'cube-array'; +const isViewDimensionCubeOrCubeArray = (viewDimension: GPUTextureViewDimension) => + viewDimension === 'cube' || viewDimension === 'cube-array'; + const s_u32 = new Uint32Array(1); const s_f32 = new Float32Array(s_u32.buffer); const s_i32 = new Int32Array(s_u32.buffer); @@ -2040,7 +2049,10 @@ const kRGBAComponents = [ const kRComponent = [TexelComponent.R] as const; -function texelsApproximatelyEqual( +/** + * Compares two Texels + */ +export function texelsApproximatelyEqual( gotRGBA: PerTexelComponent, gotFormat: GPUTextureFormat, expectRGBA: PerTexelComponent, @@ -2610,6 +2622,13 @@ function getEffectiveViewDimension( ); } +/** + * Reads a texture to an array of TexelViews, one per mip level. + * format is the format of the TexelView you want. Often this is + * same as the texture.format but if the texture.format is not + * "Encodable" then you need to choose a different format. + * Example: depth24plus -> r32float, bc1-rgba-unorm to rgba32float + */ export async function readTextureToTexelViews( t: GPUTest, texture: GPUTexture, @@ -2622,78 +2641,95 @@ export async function readTextureToTexelViews( new Map(); s_readTextureToRGBA32DeviceToPipeline.set(device, viewDimensionToPipelineMap); + const { componentType, resultType } = getTextureFormatTypeInfo(texture.format); const viewDimension = getEffectiveViewDimension(t, descriptor); - const id = `${viewDimension}:${texture.sampleCount}`; + const id = `${texture.format}:${viewDimension}:${texture.sampleCount}`; let pipeline = viewDimensionToPipelineMap.get(id); if (!pipeline) { let textureWGSL; let loadWGSL; - let dimensionWGSL = 'textureDimensions(tex, uni.mipLevel)'; + let dimensionWGSL = 'textureDimensions(tex, 0)'; switch (viewDimension) { case '2d': if (texture.sampleCount > 1) { - textureWGSL = 'texture_multisampled_2d'; + textureWGSL = `texture_multisampled_2d<${componentType}>`; loadWGSL = 'textureLoad(tex, coord.xy, sampleIndex)'; dimensionWGSL = 'textureDimensions(tex)'; } else { - textureWGSL = 'texture_2d'; - loadWGSL = 'textureLoad(tex, coord.xy, mipLevel)'; + textureWGSL = `texture_2d<${componentType}>`; + loadWGSL = 'textureLoad(tex, coord.xy, 0)'; } break; case 'cube-array': // cube-array doesn't exist in compat so we can just use 2d_array for this case '2d-array': - textureWGSL = 'texture_2d_array'; + textureWGSL = `texture_2d_array<${componentType}>`; loadWGSL = ` textureLoad( tex, coord.xy, coord.z, - mipLevel)`; + 0)`; break; case '3d': - textureWGSL = 'texture_3d'; - loadWGSL = 'textureLoad(tex, coord.xyz, mipLevel)'; + textureWGSL = `texture_3d<${componentType}>`; + loadWGSL = 'textureLoad(tex, coord.xyz, 0)'; break; case 'cube': - textureWGSL = 'texture_cube'; + textureWGSL = `texture_cube<${componentType}>`; loadWGSL = ` - textureLoadCubeAs2DArray(tex, coord.xy, coord.z, mipLevel); + textureLoadCubeAs2DArray(tex, coord.xy, coord.z); `; break; + case '1d': + textureWGSL = `texture_1d<${componentType}>`; + loadWGSL = `textureLoad(tex, coord.x, 0)`; + dimensionWGSL = `vec2u(textureDimensions(tex), 1)`; + break; default: unreachable(`unsupported view: ${viewDimension}`); } + + const textureLoadCubeWGSL = ` + const faceMat = array( + mat3x3f( 0, 0, -2, 0, -2, 0, 1, 1, 1), // pos-x + mat3x3f( 0, 0, 2, 0, -2, 0, -1, 1, -1), // neg-x + mat3x3f( 2, 0, 0, 0, 0, 2, -1, 1, -1), // pos-y + mat3x3f( 2, 0, 0, 0, 0, -2, -1, -1, 1), // neg-y + mat3x3f( 2, 0, 0, 0, -2, 0, -1, 1, 1), // pos-z + mat3x3f(-2, 0, 0, 0, -2, 0, 1, 1, -1)); // neg-z + + // needed for compat mode. + fn textureLoadCubeAs2DArray(tex: texture_cube<${componentType}>, coord: vec2u, layer: u32) -> ${resultType} { + // convert texel coord normalized coord + let size = textureDimensions(tex, 0); + let uv = (vec2f(coord) + 0.5) / vec2f(size.xy); + + // convert uv + layer into cube coord + let cubeCoord = faceMat[layer] * vec3f(uv, 1.0); + + // We have to use textureGather as it's the only texture builtin that works on cubemaps + // with integer texture formats. + let r = textureGather(0, tex, smp, cubeCoord); + let g = textureGather(1, tex, smp, cubeCoord); + let b = textureGather(2, tex, smp, cubeCoord); + let a = textureGather(3, tex, smp, cubeCoord); + + // element 3 is the texel corresponding to cubeCoord + return ${resultType}(r[3], g[3], b[3], a[3]); + } + `; + const module = device.createShaderModule({ code: ` - const faceMat = array( - mat3x3f( 0, 0, -2, 0, -2, 0, 1, 1, 1), // pos-x - mat3x3f( 0, 0, 2, 0, -2, 0, -1, 1, -1), // neg-x - mat3x3f( 2, 0, 0, 0, 0, 2, -1, 1, -1), // pos-y - mat3x3f( 2, 0, 0, 0, 0, -2, -1, -1, 1), // neg-y - mat3x3f( 2, 0, 0, 0, -2, 0, -1, 1, 1), // pos-z - mat3x3f(-2, 0, 0, 0, -2, 0, 1, 1, -1)); // neg-z - - // needed for compat mode. - fn textureLoadCubeAs2DArray(tex: texture_cube, coord: vec2u, layer: u32, mipLevel: u32) -> vec4f { - // convert texel coord normalized coord - let size = textureDimensions(tex, mipLevel); - let uv = (vec2f(coord) + 0.5) / vec2f(size.xy); - - // convert uv + layer into cube coord - let cubeCoord = faceMat[layer] * vec3f(uv, 1.0); - - return textureSampleLevel(tex, smp, cubeCoord, f32(mipLevel)); - } - + ${isViewDimensionCubeOrCubeArray(viewDimension) ? textureLoadCubeWGSL : ''} struct Uniforms { - mipLevel: u32, sampleCount: u32, }; @group(0) @binding(0) var uni: Uniforms; @group(0) @binding(1) var tex: ${textureWGSL}; @group(0) @binding(2) var smp: sampler; - @group(0) @binding(3) var data: array; + @group(0) @binding(3) var data: array<${resultType}>; @compute @workgroup_size(1) fn cs( @builtin(global_invocation_id) global_invocation_id : vec3) { @@ -2704,12 +2740,52 @@ export async function readTextureToTexelViews( global_invocation_id.x; let coord = vec3u(global_invocation_id.x / uni.sampleCount, global_invocation_id.yz); let sampleIndex = global_invocation_id.x % uni.sampleCount; - let mipLevel = uni.mipLevel; data[ndx] = ${loadWGSL}; } `, }); - pipeline = device.createComputePipeline({ layout: 'auto', compute: { module } }); + const sampleType = isDepthTextureFormat(texture.format) + ? 'unfilterable-float' + : isStencilTextureFormat(texture.format) + ? 'uint' + : kTextureFormatInfo[texture.format].color?.type ?? 'unfilterable-float'; + const bindGroupLayout = device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'uniform', + }, + }, + { + binding: 1, + visibility: GPUShaderStage.COMPUTE, + texture: { + sampleType, + viewDimension, + }, + }, + { + binding: 2, + visibility: GPUShaderStage.COMPUTE, + sampler: { + type: 'non-filtering', + }, + }, + { + binding: 3, + visibility: GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + }, + }, + ], + }); + const layout = device.createPipelineLayout({ + bindGroupLayouts: [bindGroupLayout], + }); + pipeline = device.createComputePipeline({ layout, compute: { module } }); viewDimensionToPipelineMap.set(id, pipeline); } @@ -2719,7 +2795,7 @@ export async function readTextureToTexelViews( for (let mipLevel = 0; mipLevel < texture.mipLevelCount; ++mipLevel) { const size = virtualMipSize(texture.dimension, texture, mipLevel); - const uniformValues = new Uint32Array([mipLevel, texture.sampleCount, 0, 0]); // min size is 16 bytes + const uniformValues = new Uint32Array([texture.sampleCount, 0, 0, 0]); // min size is 16 bytes const uniformBuffer = t.createBufferTracked({ size: uniformValues.byteLength, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, @@ -2749,6 +2825,8 @@ export async function readTextureToTexelViews( resource: texture.createView({ dimension: viewDimension, aspect, + baseMipLevel: mipLevel, + mipLevelCount: 1, }), }, { binding: 2, resource: sampler }, @@ -2772,7 +2850,9 @@ export async function readTextureToTexelViews( await readBuffer.mapAsync(GPUMapMode.READ); // need a copy of the data since unmapping will nullify the typedarray view. - const data = new Float32Array(readBuffer.getMappedRange()).slice(); + const Ctor = + componentType === 'i32' ? Int32Array : componentType === 'u32' ? Uint32Array : Float32Array; + const data = new Ctor(readBuffer.getMappedRange()).slice(); readBuffer.unmap(); const { sampleCount } = texture; @@ -3183,11 +3263,13 @@ async function identifySamplePoints( } const pad2 = (n: number) => n.toString().padStart(2); + const pad3 = (n: number) => n.toString().padStart(3); const fix5 = (n: number) => n.toFixed(5); + const formatValue = isSintOrUintFormat(format) ? pad3 : fix5; const formatTexel = (texel: PerTexelComponent | undefined) => texel ? Object.entries(texel) - .map(([k, v]) => `${k}: ${fix5(v)}`) + .map(([k, v]) => `${k}: ${formatValue(v)}`) .join(', ') : '*texel values unavailable*';