From c5ab6ed9d5bb23da813730ac586f8105a1b59c57 Mon Sep 17 00:00:00 2001 From: Greggman Date: Tue, 6 Aug 2024 10:19:34 -0700 Subject: [PATCH] WGSL textureSampleLevel execution tests (#3888) Cube and cube-arrays are not thoroughly tested at the moment. See TODO in comments. --- .../call/builtin/textureSample.spec.ts | 103 +-- .../call/builtin/textureSampleLevel.spec.ts | 640 ++++++++++++-- .../expression/call/builtin/texture_utils.ts | 791 +++++++++++++----- src/webgpu/util/texture.ts | 69 +- 4 files changed, 1218 insertions(+), 385 deletions(-) diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts index d43df36facb5..e1aa3f67328c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSample.spec.ts @@ -5,14 +5,12 @@ note: uniformity validation is covered in src/webgpu/shader/validation/uniformit `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { unreachable } from '../../../../../../common/util/util.js'; import { isCompressedTextureFormat, kCompressedTextureFormats, kEncodableTextureFormats, - kTextureFormatInfo, } from '../../../../../format_info.js'; -import { GPUTest, TextureTestMixin } from '../../../../../gpu_test.js'; +import { TextureTestMixin } from '../../../../../gpu_test.js'; import { vec2, @@ -29,38 +27,17 @@ import { kCubeSamplePointMethods, SamplePointMethods, chooseTextureSize, + isPotentiallyFilterableAndFillable, + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable, + getDepthOrArrayLayersForViewDimension, + getTextureTypeForTextureViewDimension, + WGSLTextureSampleTest, } from './texture_utils.js'; import { generateCoordBoundaries, generateOffsets } from './utils.js'; const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; -function getDepthOrArrayLayersForViewDimension(viewDimension: GPUTextureViewDimension) { - switch (viewDimension) { - case '2d': - return 1; - case '3d': - return 8; - case 'cube': - return 6; - default: - unreachable(); - } -} - -function getTextureTypeForTextureViewDimension(viewDimension: GPUTextureViewDimension) { - switch (viewDimension) { - case '2d': - return 'texture_2d'; - case '3d': - return 'texture_3d'; - case 'cube': - return 'texture_cube'; - default: - unreachable(); - } -} - -export const g = makeTestGroup(TextureTestMixin(GPUTest)); +export const g = makeTestGroup(TextureTestMixin(WGSLTextureSampleTest)); g.test('sampled_1d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesample') @@ -103,14 +80,7 @@ Parameters: .params(u => u .combine('format', kTestableColorFormats) - .filter(t => { - const type = kTextureFormatInfo[t.format].color?.type; - const canPotentialFilter = type === 'float' || type === 'unfilterable-float'; - // We can't easily put random bytes into compressed textures if they are float formats - // since we want the range to be +/- 1000 and not +/- infinity or NaN. - const isFillable = !isCompressedTextureFormat(t.format) || !t.format.endsWith('float'); - return canPotentialFilter && isFillable; - }) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('samplePoints', kSamplePointMethods) .beginSubcases() .combine('addressModeU', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) @@ -118,16 +88,9 @@ Parameters: .combine('minFilter', ['nearest', 'linear'] as const) .combine('offset', [false, true] as const) ) - .beforeAllSubcases(t => { - const { format } = t.params; - t.skipIfTextureFormatNotSupported(format); - const info = kTextureFormatInfo[format]; - if (info.color?.type === 'unfilterable-float') { - t.selectDeviceOrSkipTestCase('float32-filterable'); - } else { - t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); - } - }) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) + ) .fn(async t => { const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; @@ -194,14 +157,7 @@ test mip level selection based on derivatives .params(u => u .combine('format', kTestableColorFormats) - .filter(t => { - const type = kTextureFormatInfo[t.format].color?.type; - const canPotentialFilter = type === 'float' || type === 'unfilterable-float'; - // We can't easily put random bytes into compressed textures if they are float formats - // since we want the range to be +/- 1000 and not +/- infinity or NaN. - const isFillable = !isCompressedTextureFormat(t.format) || !t.format.endsWith('float'); - return canPotentialFilter && isFillable; - }) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('mipmapFilter', ['nearest', 'linear'] as const) .beginSubcases() // note: this is the derivative we want at sample time. It is not the value @@ -220,16 +176,9 @@ test mip level selection based on derivatives { ddx: 1.5, ddy: 1.5, uvwStart: [-3.5, -4] as const }, // test mix between 1 and 2 with negative coords ]) ) - .beforeAllSubcases(t => { - const { format } = t.params; - t.skipIfTextureFormatNotSupported(format); - const info = kTextureFormatInfo[format]; - if (info.color?.type === 'unfilterable-float') { - t.selectDeviceOrSkipTestCase('float32-filterable'); - } else { - t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); - } - }) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) + ) .fn(async t => { const { format, mipmapFilter, ddx, ddy, uvwStart, offset } = t.params; @@ -285,14 +234,7 @@ Parameters: .params(u => u .combine('format', kTestableColorFormats) - .filter(t => { - const type = kTextureFormatInfo[t.format].color?.type; - const canPotentialFilter = type === 'float' || type === 'unfilterable-float'; - // We can't easily put random bytes into compressed textures if they are float formats - // since we want the range to be +/- 1000 and not +/- infinity or NaN. - const isFillable = !isCompressedTextureFormat(t.format) || !t.format.endsWith('float'); - return canPotentialFilter && isFillable; - }) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('viewDimension', ['3d', 'cube'] as const) .filter(t => !isCompressedTextureFormat(t.format) || t.viewDimension === 'cube') .combine('samplePoints', kCubeSamplePointMethods) @@ -305,16 +247,9 @@ Parameters: .combine('offset', [false, true] as const) .filter(t => t.viewDimension !== 'cube' || t.offset !== true) ) - .beforeAllSubcases(t => { - const { format } = t.params; - t.skipIfTextureFormatNotSupported(format); - const info = kTextureFormatInfo[format]; - if (info.color?.type === 'unfilterable-float') { - t.selectDeviceOrSkipTestCase('float32-filterable'); - } else { - t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); - } - }) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) + ) .fn(async t => { const { format, diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts index f8073c65d66d..729563553260 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts @@ -3,14 +3,54 @@ Samples a texture. Must only be used in a fragment shader stage. Must only be invoked in uniform control flow. + +- TODO: Test un-encodable formats. +- TODO: set mipLevelCount to 3 for cubemaps. See MAINTENANCE_TODO below + + The issue is sampling a corner of a cubemap is undefined. We try to quantize coordinates + so we never get a corner but when sampling smaller mip levels that's more difficult unless we make the textures + larger. Larger is slower. + + Solution 1: Fix the quantization + Solution 2: special case checking cube corners. Expect some value between the color of the 3 corner texels. + `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; +import { + isCompressedTextureFormat, + isDepthTextureFormat, + isEncodableTextureFormat, + kCompressedTextureFormats, + kDepthStencilFormats, + kEncodableTextureFormats, +} from '../../../../../format_info.js'; + +import { + appendComponentTypeForFormatToTextureType, + checkCallResults, + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + doTextureCalls, + generateSamplePointsCube, + generateTextureBuiltinInputs2D, + generateTextureBuiltinInputs3D, + getDepthOrArrayLayersForViewDimension, + getTextureTypeForTextureViewDimension, + isPotentiallyFilterableAndFillable, + kCubeSamplePointMethods, + kSamplePointMethods, + SamplePointMethods, + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable, + TextureCall, + vec2, + vec3, + WGSLTextureSampleTest, +} from './texture_utils.js'; -import { generateCoordBoundaries, generateOffsets } from './utils.js'; +const kTestableColorFormats = [...kEncodableTextureFormats, ...kCompressedTextureFormats] as const; -export const g = makeTestGroup(GPUTest); +export const g = makeTestGroup(WGSLTextureSampleTest); g.test('sampled_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -36,14 +76,70 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('coords', generateCoordBoundaries(2)) - .combine('offset', generateOffsets(2)) - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('format', kTestableColorFormats) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('addressModeU', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('addressModeV', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) + .combine('offset', [false, true] as const) + ) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) - .unimplemented(); + .fn(async t => { + const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; + + // We want at least 4 blocks or something wide enough for 3 mip levels. + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height }, + mipLevelCount: 3, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU, + addressModeV, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: 'f32' }, + offset, + hashInputs: [format, samplePoints, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, mipLevel, offset }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: 'f', + offset, + }; + }); + const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format); + const viewDescriptor = {}; + 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('sampled_array_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -51,8 +147,8 @@ g.test('sampled_array_2d_coords') ` C is i32 or u32 -fn textureSampleLevel(t: texture_2d_array, s: sampler, coords: vec2, array_index: C, level: f32) -> vec4 -fn textureSampleLevel(t: texture_2d_array, s: sampler, coords: vec2, array_index: C, level: f32, offset: vec2) -> vec4 +fn textureSampleLevel(t: texture_2d_array, s: sampler, coords: vec2, array_index: A, level: f32) -> vec4 +fn textureSampleLevel(t: texture_2d_array, s: sampler, coords: vec2, array_index: A, level: f32, offset: vec2) -> vec4 Parameters: * t The sampled or depth texture to sample. @@ -72,17 +168,76 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(2)) - .combine('offset', generateOffsets(2)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('format', kTestableColorFormats) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('A', ['i32', 'u32'] as const) + .combine('addressModeU', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('addressModeV', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) + .combine('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) + ) + .fn(async t => { + const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params; + + // We want at least 4 blocks or something wide enough for 3 mip levels. + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const depthOrArrayLayers = 4; + + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height, depthOrArrayLayers }, + mipLevelCount: 3, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU, + addressModeV, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: 'f32' }, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + offset, + hashInputs: [format, samplePoints, A, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, mipLevel, arrayIndex, offset }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: 'f', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + offset, + }; + }); + const textureType = appendComponentTypeForFormatToTextureType('texture_2d_array', format); + const viewDescriptor = {}; + 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('sampled_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -111,22 +266,96 @@ Parameters: ) .params(u => u - .combine('texture_type', ['texture_3d', 'texture_cube'] as const) + .combine('format', kTestableColorFormats) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) + .combine('viewDimension', ['3d', 'cube'] as const) + .filter(t => !isCompressedTextureFormat(t.format) || t.viewDimension === 'cube') .beginSubcases() - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('coords', generateCoordBoundaries(3)) - .combine('offset', generateOffsets(3)) - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('samplePoints', kCubeSamplePointMethods) + .filter(t => t.samplePoints !== 'cube-edges' || t.viewDimension !== '3d') + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) + .combine('offset', [false, true] as const) + .filter(t => t.viewDimension !== 'cube' || t.offset !== true) + ) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) - .unimplemented(); + .fn(async t => { + const { format, viewDimension, samplePoints, addressMode, minFilter, offset } = t.params; + + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); + + const descriptor: GPUTextureDescriptor = { + format, + dimension: viewDimension === '3d' ? '3d' : '2d', + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size: { width, height, depthOrArrayLayers }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + // MAINTENANCE_TODO: make mipLevelCount always 3 + mipLevelCount: viewDimension === 'cube' ? 1 : 3, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + addressModeW: addressMode, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = ( + viewDimension === '3d' + ? generateTextureBuiltinInputs3D(50, { + method: samplePoints as SamplePointMethods, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: 'f32' }, + offset, + hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter, offset], + }) + : generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: 'f32' }, + hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter, offset], + }) + ).map(({ coords, mipLevel, offset }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: 'f', + offset, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = getTextureTypeForTextureViewDimension(viewDimension); + 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('sampled_array_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') .desc( ` -C is i32 or u32 +A is i32 or u32 -fn textureSampleLevel(t: texture_cube_array, s: sampler, coords: vec3, array_index: C, level: f32) -> vec4 +fn textureSampleLevel(t: texture_cube_array, s: sampler, coords: vec3, array_index: A, level: f32) -> vec4 Parameters: * t The sampled or depth texture to sample. @@ -138,25 +367,90 @@ Parameters: * For the functions where level is a f32, fractional values may interpolate between two levels if the format is filterable according to the Texture Format Capabilities. * When not specified, mip level 0 is sampled. - * offset - * The optional texel offset applied to the unnormalized texture coordinate before sampling the texture. - * This offset is applied before applying any texture wrapping modes. - * The offset expression must be a creation-time expression (e.g. vec2(1, 2)). - * Each offset component must be at least -8 and at most 7. - Values outside of this range will result in a shader-creation error. + +- TODO: set mipLevelCount to 3 for cubemaps. See MAINTENANCE_TODO below + + The issue is sampling a corner of a cubemap is undefined. We try to quantize coordinates + so we never get a corner but when sampling smaller mip levels that's more difficult. + + * Solution 1: Fix the quantization + * Solution 2: special case checking cube corners. Expect some value between the color of the 3 corner texels. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(3)) - .combine('offset', generateOffsets(3)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('format', kTestableColorFormats) + .filter(t => isPotentiallyFilterableAndFillable(t.format)) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('A', ['i32', 'u32'] as const) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format); + t.skipIfTextureViewDimensionNotSupported('cube-array'); + }) + .fn(async t => { + const { format, samplePoints, A, addressMode, minFilter } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube-array'; + const size = chooseTextureSize({ + minSize: 8, + minBlocks: 4, + format, + viewDimension, + }); + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + // MAINTENANCE_TODO: Set this to 3. See above. + mipLevelCount: 1, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + addressModeW: addressMode, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: 'f32' }, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter], + }).map(({ coords, mipLevel, arrayIndex }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: 'f', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = getTextureTypeForTextureViewDimension(viewDimension); + 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('depth_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -164,8 +458,8 @@ g.test('depth_2d_coords') ` C is i32 or u32 -fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2, level: C) -> f32 -fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2, level: C, offset: vec2) -> f32 +fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2, level: L) -> f32 +fn textureSampleLevel(t: texture_depth_2d, s: sampler, coords: vec2, level: L, offset: vec2) -> f32 Parameters: * t The sampled or depth texture to sample. @@ -184,16 +478,76 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(2)) - .combine('offset', generateOffsets(2)) - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('format', kDepthStencilFormats) + // filter out stencil only formats + .filter(t => isDepthTextureFormat(t.format)) + // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. + .filter(t => isEncodableTextureFormat(t.format)) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) + .combine('L', ['i32', 'u32'] as const) + .combine('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) + ) + .fn(async t => { + const { format, samplePoints, addressMode, minFilter, L, offset } = t.params; + + // We want at least 4 blocks or something wide enough for 3 mip levels. + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height }, + mipLevelCount: 3, + usage: + GPUTextureUsage.COPY_DST | + GPUTextureUsage.TEXTURE_BINDING | + GPUTextureUsage.RENDER_ATTACHMENT, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: L }, + offset, + hashInputs: [format, samplePoints, addressMode, minFilter, L, offset], + }).map(({ coords, mipLevel, offset }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: L === 'i32' ? 'i' : 'u', + offset, + }; + }); + const textureType = appendComponentTypeForFormatToTextureType('texture_depth_2d', format); + const viewDescriptor = {}; + 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('depth_array_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -201,8 +555,8 @@ g.test('depth_array_2d_coords') ` C is i32 or u32 -fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: C, level: C) -> f32 -fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: C, level: C, offset: vec2) -> f32 +fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: A, level: L) -> f32 +fn textureSampleLevel(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: A, level: L, offset: vec2) -> f32 Parameters: * t The sampled or depth texture to sample. @@ -222,17 +576,81 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(2)) - .combine('offset', generateOffsets(2)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('format', kDepthStencilFormats) + // filter out stencil only formats + .filter(t => isDepthTextureFormat(t.format)) + // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. + .filter(t => isEncodableTextureFormat(t.format)) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) + .combine('A', ['i32', 'u32'] as const) + .combine('L', ['i32', 'u32'] as const) + .combine('offset', [false, true] as const) + ) + .beforeAllSubcases(t => + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) - .unimplemented(); + .fn(async t => { + const { format, samplePoints, addressMode, minFilter, A, L, offset } = t.params; + + // We want at least 4 blocks or something wide enough for 3 mip levels. + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height }, + mipLevelCount: 3, + usage: + GPUTextureUsage.COPY_DST | + GPUTextureUsage.TEXTURE_BINDING | + GPUTextureUsage.RENDER_ATTACHMENT, + ...(t.isCompatibility && { textureBindingViewDimension: '2d-array' }), + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + mipLevel: { num: texture.mipLevelCount, type: L }, + offset, + hashInputs: [format, samplePoints, addressMode, minFilter, L, A, offset], + }).map(({ coords, mipLevel, arrayIndex, offset }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: L === 'i32' ? 'i' : 'u', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + offset, + }; + }); + const textureType = appendComponentTypeForFormatToTextureType('texture_depth_2d_array', format); + const viewDescriptor: GPUTextureViewDescriptor = { dimension: '2d-array' }; + 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('depth_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') @@ -240,8 +658,8 @@ g.test('depth_3d_coords') ` C is i32 or u32 -fn textureSampleLevel(t: texture_depth_cube, s: sampler, coords: vec3, level: C) -> f32 -fn textureSampleLevel(t: texture_depth_cube_array, s: sampler, coords: vec3, array_index: C, level: C) -> f32 +fn textureSampleLevel(t: texture_depth_cube, s: sampler, coords: vec3, level: L) -> f32 +fn textureSampleLevel(t: texture_depth_cube_array, s: sampler, coords: vec3, array_index: A, level: L) -> f32 Parameters: * t The sampled or depth texture to sample. @@ -262,13 +680,87 @@ Parameters: ) .params(u => u - .combine('texture_type', ['texture_depth_cube', 'texture_depth_cube_array'] as const) + .combine('format', kDepthStencilFormats) + // filter out stencil only formats + .filter(t => isDepthTextureFormat(t.format)) + // MAINTENANCE_TODO: Remove when support for depth24plus, depth24plus-stencil8, and depth32float-stencil8 is added. + .filter(t => isEncodableTextureFormat(t.format)) + .combineWithParams([ + { viewDimension: 'cube' }, + { viewDimension: 'cube-array', A: 'i32' }, + { viewDimension: 'cube-array', A: 'u32' }, + ] as const) .beginSubcases() - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(3)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('level', [undefined, 0, 1, 'textureNumLevels', 'textureNumLevels+1'] as const) + .combine('samplePoints', kCubeSamplePointMethods) + .combine('L', ['i32', 'u32'] as const) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('minFilter', ['nearest', 'linear'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format); + t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); + }) + .fn(async t => { + const { format, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params; + + const size = chooseTextureSize({ + minSize: 8, + minBlocks: 4, + format, + viewDimension, + }); + const descriptor: GPUTextureDescriptor = { + format, + size, + usage: + GPUTextureUsage.COPY_DST | + GPUTextureUsage.TEXTURE_BINDING | + GPUTextureUsage.RENDER_ATTACHMENT, + mipLevelCount: 3, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + addressModeW: addressMode, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + mipLevel: { num: texture.mipLevelCount, type: L }, + arrayIndex: A ? { num: texture.depthOrArrayLayers, type: A } : undefined, + hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter], + }).map(({ coords, mipLevel, arrayIndex }) => { + return { + builtin: 'textureSampleLevel', + coordType: 'f', + coords, + mipLevel, + levelType: L === 'i32' ? 'i' : 'u', + arrayIndex, + arrayIndexType: A ? (A === 'i32' ? 'i' : 'u') : undefined, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = + viewDimension === 'cube' ? 'texture_depth_cube' : 'texture_depth_cube_array'; + 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); + }); 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 674d09bb1f1a..8177f51ef30c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -8,8 +8,11 @@ import { kEncodableTextureFormats, kTextureFormatInfo, } from '../../../../../format_info.js'; -import { GPUTest, TextureTestMixinType } from '../../../../../gpu_test.js'; -import { float32ToUint32 } from '../../../../../util/conversion.js'; +import { + GPUTest, + GPUTestSubcaseBatchState, + TextureTestMixinType, +} from '../../../../../gpu_test.js'; import { align, clamp, @@ -52,6 +55,209 @@ export const kSampleTypeInfo = { }, } as const; +/** + * Return the texture type for a given view dimension + */ +export function getTextureTypeForTextureViewDimension(viewDimension: GPUTextureViewDimension) { + switch (viewDimension) { + case '1d': + return 'texture_1d'; + case '2d': + return 'texture_2d'; + case '2d-array': + return 'texture_2d_array'; + case '3d': + return 'texture_3d'; + case 'cube': + return 'texture_cube'; + case 'cube-array': + return 'texture_cube_array'; + default: + unreachable(); + } +} + +/** + * Returns if a texture format can potentially be filtered and can be filled with random data. + */ +export function isPotentiallyFilterableAndFillable(format: GPUTextureFormat) { + const type = kTextureFormatInfo[format].color?.type; + const canPotentiallyFilter = type === 'float' || type === 'unfilterable-float'; + // We can't easily put random bytes into compressed textures if they are float formats + // since we want the range to be +/- 1000 and not +/- infinity or NaN. + const isFillable = !isCompressedTextureFormat(format) || !format.endsWith('float'); + return canPotentiallyFilter && isFillable; +} + +/** + * skips the test if the texture format is not supported or not available or not filterable. + */ +export function skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable( + t: GPUTestSubcaseBatchState, + format: GPUTextureFormat +) { + t.skipIfTextureFormatNotSupported(format); + const info = kTextureFormatInfo[format]; + if (info.color?.type === 'unfilterable-float') { + t.selectDeviceOrSkipTestCase('float32-filterable'); + } else { + t.selectDeviceForTextureFormatOrSkipTestCase(format); + } +} + +/** + * Gets the mip gradient values for the current device. + * The issue is, different GPUs have different ways of mixing between mip levels. + * For most GPUs it's linear but for AMD GPUs on Mac in particular, it's something + * else (which AFAICT is against all the specs). + * + * We seemingly have 3 options: + * + * 1. Increase the tolerances of tests so they pass on AMD. + * 2. Mark AMD as failing + * 3. Try to figure out how the GPU converts mip levels into weights + * + * We're doing 3. + * + * There's an assumption that the gradient will be the same for all formats + * and usages. + */ +const kMipGradientSteps = 16; +const s_deviceToMipGradientValues = new WeakMap(); +async function initMipGradientValuesForDevice(t: GPUTest) { + const { device } = t; + const weights = s_deviceToMipGradientValues.get(device); + if (!weights) { + const module = device.createShaderModule({ + code: ` + @group(0) @binding(0) var tex: texture_2d; + @group(0) @binding(1) var smp: sampler; + @group(0) @binding(2) var result: array; + + @compute @workgroup_size(1) fn cs(@builtin(global_invocation_id) id: vec3u) { + let mipLevel = f32(id.x) / ${kMipGradientSteps}; + result[id.x] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r; + } + `, + }); + + const pipeline = device.createComputePipeline({ + layout: 'auto', + compute: { module }, + }); + + const texture = t.createTextureTracked({ + size: [2, 2, 1], + format: 'r8unorm', + usage: GPUTextureUsage.TEXTURE_BINDING | GPUTextureUsage.COPY_DST, + mipLevelCount: 2, + }); + + device.queue.writeTexture( + { texture, mipLevel: 1 }, + new Uint8Array([255]), + { bytesPerRow: 1 }, + [1, 1] + ); + + const sampler = device.createSampler({ + minFilter: 'linear', + magFilter: 'linear', + mipmapFilter: 'linear', + }); + + const storageBuffer = t.createBufferTracked({ + size: 4 * (kMipGradientSteps + 1), + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const resultBuffer = t.createBufferTracked({ + size: storageBuffer.size, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, + }); + + const bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: texture.createView() }, + { binding: 1, resource: sampler }, + { binding: 2, resource: { buffer: storageBuffer } }, + ], + }); + + const encoder = device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(kMipGradientSteps + 1); + pass.end(); + encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size); + device.queue.submit([encoder.finish()]); + + await resultBuffer.mapAsync(GPUMapMode.READ); + const weights = Array.from(new Float32Array(resultBuffer.getMappedRange())); + resultBuffer.unmap(); + + texture.destroy(); + storageBuffer.destroy(); + resultBuffer.destroy(); + + // Validate the weights + assert(weights[0] === 0); + assert(weights[kMipGradientSteps] === 1); + assert(weights[kMipGradientSteps / 2] === 0.5); + + // Note: for 16 steps, these are the AMD weights + // + // standard + // step mipLevel gpu AMD + // ---- -------- -------- ---------- + // 0: 0 0 1 + // 1: 0.0625 0.0625 0 + // 2: 0.125 0.125 0.03125 + // 3: 0.1875 0.1875 0.109375 + // 4: 0.25 0.25 0.1875 + // 5: 0.3125 0.3125 0.265625 + // 6: 0.375 0.375 0.34375 + // 7: 0.4375 0.4375 0.421875 + // 8: 0.5 0.5 0.5 + // 9: 0.5625 0.5625 0.578125 + // 10: 0.625 0.625 0.65625 + // 11: 0.6875 0.6875 0.734375 + // 12: 0.75 0.75 0.8125 + // 13: 0.8125 0.8125 0.890625 + // 14: 0.875 0.875 0.96875 + // 15: 0.9375 0.9375 1 + // 16: 1 1 1 + // + // notice step 1 is 0 and step 15 is 1. + // so we only check the 1 through 14. + for (let i = 1; i < kMipGradientSteps - 1; ++i) { + assert(weights[i] < weights[i + 1]); + } + + s_deviceToMipGradientValues.set(device, weights); + } +} + +function getWeightForMipLevel(t: GPUTest, mipLevelCount: number, mipLevel: number) { + if (mipLevel < 0 || mipLevel >= mipLevelCount) { + return 1; + } + // linear interpolate between weights + const weights = s_deviceToMipGradientValues.get(t.device); + assert( + !!weights, + 'you must use WGSLTextureSampleTest or call initializeDeviceMipWeights before calling this function' + ); + const steps = weights.length - 1; + const w = (mipLevel % 1) * steps; + const lowerNdx = Math.floor(w); + const upperNdx = Math.ceil(w); + const mix = w % 1; + return lerp(weights[lowerNdx], weights[upperNdx], mix); +} + /** * Used for textureDimension, textureNumLevels, textureNumLayers */ @@ -94,26 +300,24 @@ export class WGSLTextureQueryTest extends GPUTest { } /** - * Generates an array of pseudo random values based on a hash. - * For `i32` generates an integer in the range [-1, num] - * For `u32` generates an integer in the range [0, num) - * for `f32` generates an number in the range [-1 to num) + * Used for textureSampleXXX */ -export function makeRepeatableValuesInRanges({ - hashInputs, - rangeDefs, -}: { - hashInputs: (number | string)[]; - rangeDefs: RangeDef[]; -}): number[] { - const _hashInputs = hashInputs.map(v => (typeof v === 'string' ? sumOfCharCodesOfString(v) : v)); - return rangeDefs.map(({ num, type }, i) => { - const range = num + type === 'u32' ? 1 : 2; - const number = (hashU32(..._hashInputs, i) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1); - return type === 'f32' ? number : Math.floor(number); - }); +export class WGSLTextureSampleTest extends GPUTest { + override async init(): Promise { + await super.init(); + await initMipGradientValuesForDevice(this); + } } +/** + * Used to specify a range from [0, num) + * The type is used to determine if values should be integers and if they can be negative. + */ +export type RangeDef = { + num: number; + type: 'f32' | 'i32' | 'u32'; +}; + function getLimitValue(v: number) { switch (v) { case Number.POSITIVE_INFINITY: @@ -252,7 +456,7 @@ export type vec4 = [number, number, number, number]; export type Dimensionality = vec1 | vec2 | vec3; type TextureCallArgKeys = keyof TextureCallArgs; -const kTextureCallArgNames: TextureCallArgKeys[] = [ +const kTextureCallArgNames: readonly TextureCallArgKeys[] = [ 'coords', 'arrayIndex', 'sampleIndex', @@ -260,7 +464,7 @@ const kTextureCallArgNames: TextureCallArgKeys[] = [ 'ddx', 'ddy', 'offset', -]; +] as const; export interface TextureCallArgs { coords?: T; @@ -273,13 +477,58 @@ export interface TextureCallArgs { } export interface TextureCall extends TextureCallArgs { - builtin: 'textureSample' | 'textureLoad'; + builtin: 'textureSample' | 'textureLoad' | 'textureSampleLevel'; coordType: 'f' | 'i' | 'u'; - levelType?: 'i' | 'u'; + levelType?: 'i' | 'u' | 'f'; arrayIndexType?: 'i' | 'u'; sampleIndexType?: 'i' | 'u'; } +const s_u32 = new Uint32Array(1); +const s_f32 = new Float32Array(s_u32.buffer); +const s_i32 = new Int32Array(s_u32.buffer); + +const kBitCastFunctions = { + f: (v: number) => { + s_f32[0] = v; + return s_u32[0]; + }, + i: (v: number) => { + s_i32[0] = v; + assert(s_i32[0] === v, 'check we are not casting non-int or out-of-range value'); + return s_u32[0]; + }, + u: (v: number) => { + s_u32[0] = v; + assert(s_u32[0] === v, 'check we are not casting non-uint or out-of-range value'); + return s_u32[0]; + }, +}; + +function getCallArgType( + call: TextureCall, + argName: (typeof kTextureCallArgNames)[number] +) { + switch (argName) { + case 'coords': + return call.coordType; + case 'mipLevel': + assert(call.levelType !== undefined); + return call.levelType; + case 'arrayIndex': + assert(call.arrayIndexType !== undefined); + return call.arrayIndexType; + case 'sampleIndex': + assert(call.sampleIndexType !== undefined); + return call.sampleIndexType; + case 'ddx': + case 'ddy': + return 'f'; + default: + unreachable(); + } +} + function toArray(coords: Dimensionality): number[] { if (coords instanceof Array) { return coords; @@ -395,6 +644,7 @@ export function softwareTextureReadMipLevel( sampler: GPUSamplerDescriptor | undefined, mipLevel: number ): PerTexelComponent { + assert(mipLevel % 1 === 0); const { format } = texture.texels[0]; const rep = kTexelRepresentationInfo[format]; const textureSize = virtualMipSize( @@ -408,18 +658,31 @@ export function softwareTextureReadMipLevel( sampler?.addressModeW ?? 'clamp-to-edge', ]; - const load = (at: number[]) => - texture.texels[mipLevel].color({ + const isCube = + texture.viewDescriptor.dimension === 'cube' || + texture.viewDescriptor.dimension === 'cube-array'; + + const arrayIndexMult = isCube ? 6 : 1; + const numLayers = textureSize[2] / arrayIndexMult; + assert(numLayers % 1 === 0); + const textureSizeForCube = [textureSize[0], textureSize[1], 6]; + + const load = (at: number[]) => { + const zFromArrayIndex = + call.arrayIndex !== undefined + ? clamp(call.arrayIndex, { min: 0, max: numLayers - 1 }) * arrayIndexMult + : 0; + return texture.texels[mipLevel].color({ x: Math.floor(at[0]), y: Math.floor(at[1] ?? 0), - z: call.arrayIndex ?? Math.floor(at[2] ?? 0), + z: Math.floor(at[2] ?? 0) + zFromArrayIndex, sampleIndex: call.sampleIndex, }); - - const isCube = texture.viewDescriptor.dimension === 'cube'; + }; switch (call.builtin) { - case 'textureSample': { + case 'textureSample': + case 'textureSampleLevel': { let coords = toArray(call.coords!); if (isCube) { @@ -436,7 +699,7 @@ export function softwareTextureReadMipLevel( // ├───┼───┼───┼───┤ // │ │ │ │ b │ // └───┴───┴───┴───┘ - let at = coords.map((v, i) => v * textureSize[i] - 0.5); + let at = coords.map((v, i) => v * (isCube ? textureSizeForCube : textureSize)[i] - 0.5); // Apply offset in whole texel units // This means the offset is added at each mip level in texels. There's no @@ -571,40 +834,22 @@ export function softwareTextureReadMipLevel( } /** - * The software version of a texture builtin (eg: textureSample) - * Note that this is not a complete implementation. Rather it's only - * what's needed to generate the correct expected value for the tests. + * Reads a texture, optionally sampling between 2 mipLevels */ -export function softwareTextureRead( +export function softwareTextureReadLevel( + t: GPUTest, call: TextureCall, texture: Texture, - sampler: GPUSamplerDescriptor + sampler: GPUSamplerDescriptor | undefined, + mipLevel: number ): PerTexelComponent { - assert(call.ddx !== undefined); - assert(call.ddy !== undefined); - const rep = kTexelRepresentationInfo[texture.texels[0].format]; - const texSize = reifyExtent3D(texture.descriptor.size); - const textureSize = [texSize.width, texSize.height]; - - // ddx and ddy are the values that would be passed to textureSampleGrad - // If we're emulating textureSample then they're the computed derivatives - // such that if we passed them to textureSampleGrad they'd produce the - // same result. - const ddx: readonly number[] = typeof call.ddx === 'number' ? [call.ddx] : call.ddx; - const ddy: readonly number[] = typeof call.ddy === 'number' ? [call.ddy] : call.ddy; - - // Compute the mip level the same way textureSampleGrad does - const scaledDdx = ddx.map((v, i) => v * textureSize[i]); - const scaledDdy = ddy.map((v, i) => v * textureSize[i]); - const dotDDX = dotProduct(scaledDdx, scaledDdx); - const dotDDY = dotProduct(scaledDdy, scaledDdy); - const deltaMax = Math.max(dotDDX, dotDDY); - // MAINTENANCE_TODO: handle texture view baseMipLevel and mipLevelCount? - const mipLevel = 0.5 * Math.log2(deltaMax); - const mipLevelCount = texture.texels.length; const maxLevel = mipLevelCount - 1; + if (!sampler) { + return softwareTextureReadMipLevel(call, texture, sampler, mipLevel); + } + switch (sampler.mipmapFilter) { case 'linear': { const clampedMipLevel = clamp(mipLevel, { min: 0, max: maxLevel }); @@ -612,14 +857,14 @@ export function softwareTextureRead( const nextMipLevel = Math.ceil(clampedMipLevel); const t0 = softwareTextureReadMipLevel(call, texture, sampler, baseMipLevel); const t1 = softwareTextureReadMipLevel(call, texture, sampler, nextMipLevel); - const mix = mipLevel % 1; + const mix = getWeightForMipLevel(t, mipLevelCount, mipLevel); const values = [ { v: t0, weight: 1 - mix }, { v: t1, weight: mix }, ]; const out: PerTexelComponent = {}; for (const { v, weight } of values) { - for (const component of rep.componentOrder) { + for (const component of kRGBAComponents) { out[component] = (out[component] ?? 0) + v[component]! * weight; } } @@ -634,6 +879,40 @@ export function softwareTextureRead( } } +/** + * The software version of a texture builtin (eg: textureSample) + * Note that this is not a complete implementation. Rather it's only + * what's needed to generate the correct expected value for the tests. + */ +export function softwareTextureRead( + t: GPUTest, + call: TextureCall, + texture: Texture, + sampler: GPUSamplerDescriptor +): PerTexelComponent { + assert(call.ddx !== undefined); + assert(call.ddy !== undefined); + const texSize = reifyExtent3D(texture.descriptor.size); + const textureSize = [texSize.width, texSize.height]; + + // ddx and ddy are the values that would be passed to textureSampleGrad + // If we're emulating textureSample then they're the computed derivatives + // such that if we passed them to textureSampleGrad they'd produce the + // same result. + const ddx: readonly number[] = typeof call.ddx === 'number' ? [call.ddx] : call.ddx; + const ddy: readonly number[] = typeof call.ddy === 'number' ? [call.ddy] : call.ddy; + + // Compute the mip level the same way textureSampleGrad does + const scaledDdx = ddx.map((v, i) => v * textureSize[i]); + const scaledDdy = ddy.map((v, i) => v * textureSize[i]); + const dotDDX = dotProduct(scaledDdx, scaledDdx); + const dotDDY = dotProduct(scaledDdy, scaledDdy); + const deltaMax = Math.max(dotDDX, dotDDY); + // MAINTENANCE_TODO: handle texture view baseMipLevel and mipLevelCount? + const mipLevel = 0.5 * Math.log2(deltaMax); + return softwareTextureReadLevel(t, call, texture, sampler, mipLevel); +} + export type TextureTestOptions = { ddx?: number; // the derivative we want at sample time ddy?: number; @@ -651,10 +930,8 @@ export type TextureTestOptions = { */ function isOutOfBoundsCall(texture: Texture, call: TextureCall) { assert(call.coords !== undefined); - assert(call.offset === undefined); const desc = reifyTextureDescriptor(texture.descriptor); - const { coords, mipLevel, arrayIndex, sampleIndex } = call; if (mipLevel !== undefined && (mipLevel < 0 || mipLevel >= desc.mipLevelCount)) { @@ -690,24 +967,17 @@ function isOutOfBoundsCall(texture: Texture, call: Tex return false; } -/** - * For a texture builtin with no sampler (eg textureLoad), - * any out of bounds access is allowed to return one of: - * - * * the value of any texel in the texture - * * 0,0,0,0 or 0,0,0,1 if not a depth texture - * * 0 if a depth texture - */ -function okBecauseOutOfBounds( +function isValidOutOfBoundsValue( texture: Texture, - call: TextureCall, gotRGBA: PerTexelComponent, maxFractionalDiff: number ) { - if (!isOutOfBoundsCall(texture, call)) { - return false; - } - + // For a texture builtin with no sampler (eg textureLoad), + // any out of bounds access is allowed to return one of: + // + // * the value of any texel in the texture + // * 0,0,0,0 or 0,0,0,1 if not a depth texture + // * 0 if a depth texture if (texture.descriptor.format.includes('depth')) { if (gotRGBA.R === 0) { return true; @@ -723,6 +993,7 @@ function okBecauseOutOfBounds( } } + // Can be any texel value for (let mipLevel = 0; mipLevel < texture.texels.length; ++mipLevel) { const mipTexels = texture.texels[mipLevel]; const size = virtualMipSize( @@ -749,6 +1020,27 @@ function okBecauseOutOfBounds( return false; } +/** + * For a texture builtin with no sampler (eg textureLoad), + * any out of bounds access is allowed to return one of: + * + * * the value of any texel in the texture + * * 0,0,0,0 or 0,0,0,1 if not a depth texture + * * 0 if a depth texture + */ +function okBecauseOutOfBounds( + texture: Texture, + call: TextureCall, + gotRGBA: PerTexelComponent, + maxFractionalDiff: number +) { + if (!isOutOfBoundsCall(texture, call)) { + return false; + } + + return isValidOutOfBoundsValue(texture, gotRGBA, maxFractionalDiff); +} + const kRGBAComponents = [ TexelComponent.R, TexelComponent.G, @@ -804,7 +1096,9 @@ export async function checkCallResults( results: PerTexelComponent[] ) { const errs: string[] = []; - const rep = kTexelRepresentationInfo[texture.texels[0].format]; + const format = texture.texels[0].format; + const rep = kTexelRepresentationInfo[format]; + const size = reifyExtent3D(texture.descriptor.size); const maxFractionalDiff = sampler?.minFilter === 'linear' || sampler?.magFilter === 'linear' || @@ -812,14 +1106,12 @@ export async function checkCallResults( ? getMaxFractionalDiffForTextureFormat(texture.descriptor.format) : 0; - for (let callIdx = 0; callIdx < calls.length; callIdx++) { + for (let callIdx = 0; callIdx < calls.length && errs.length === 0; callIdx++) { const call = calls[callIdx]; const gotRGBA = results[callIdx]; - const expectRGBA = softwareTextureReadMipLevel(call, texture, sampler, call.mipLevel ?? 0); + const expectRGBA = softwareTextureReadLevel(t, call, texture, sampler, call.mipLevel ?? 0); - if ( - texelsApproximatelyEqual(gotRGBA, expectRGBA, texture.texels[0].format, maxFractionalDiff) - ) { + if (texelsApproximatelyEqual(gotRGBA, expectRGBA, format, maxFractionalDiff)) { continue; } @@ -827,8 +1119,8 @@ export async function checkCallResults( continue; } - const got = convertResultFormatToTexelViewFormat(gotRGBA, texture.texels[0].format); - const expect = convertResultFormatToTexelViewFormat(expectRGBA, texture.texels[0].format); + const got = convertResultFormatToTexelViewFormat(gotRGBA, format); + const expect = convertResultFormatToTexelViewFormat(expectRGBA, format); const gULP = rep.bitsToULPFromZero(rep.numberToBits(got)); const eULP = rep.bitsToULPFromZero(rep.numberToBits(expect)); for (const component of rep.componentOrder) { @@ -839,7 +1131,6 @@ export async function checkCallResults( const relDiff = absDiff / Math.max(Math.abs(g), Math.abs(e)); if (ulpDiff > 3 && absDiff > maxFractionalDiff) { const desc = describeTextureCall(call); - const size = reifyExtent3D(texture.descriptor.size); errs.push(`component was not as expected: size: [${size.width}, ${size.height}, ${size.depthOrArrayLayers}] mipCount: ${texture.descriptor.mipLevelCount ?? 1} @@ -854,25 +1145,26 @@ export async function checkCallResults( if (sampler) { const expectedSamplePoints = [ 'expected:', - ...(await identifySamplePoints(texture, (texels: TexelView) => { + ...(await identifySamplePoints(texture, (texels: TexelView[]) => { return Promise.resolve( - softwareTextureReadMipLevel( + softwareTextureReadLevel( + t, call, { - texels: [texels], + texels, descriptor: texture.descriptor, viewDescriptor: texture.viewDescriptor, }, sampler, - 0 + call.mipLevel ?? 0 ) ); })), ]; const gotSamplePoints = [ 'got:', - ...(await identifySamplePoints(texture, async (texels: TexelView) => { - const gpuTexture = createTextureFromTexelViews(t, [texels], texture.descriptor); + ...(await identifySamplePoints(texture, async (texels: TexelView[]) => { + const gpuTexture = createTextureFromTexelViews(t, texels, texture.descriptor); const result = ( await doTextureCalls(t, gpuTexture, texture.viewDescriptor, textureType, sampler, [ call, @@ -898,6 +1190,7 @@ export async function checkCallResults( * sampling from the given Texture. */ export function softwareRasterize( + t: GPUTest, texture: Texture, sampler: GPUSamplerDescriptor, targetSize: [number, number], @@ -955,7 +1248,7 @@ export function softwareRasterize( ddy: [0, ddy / textureSize.height] as T, offset: options.offset as T, }; - const sample = softwareTextureRead(call, texture, sampler); + const sample = softwareTextureRead(t, call, texture, sampler); const rgba = { R: 0, G: 0, B: 0, A: 1, ...sample }; const asRgba32Float = new Float32Array(rep.pack(rgba)); expData.set(asRgba32Float, (y * width + x) * 4); @@ -1099,7 +1392,9 @@ function getMaxFractionalDiffForTextureFormat(format: GPUTextureFormat) { // MAINTENANCE_TODO: Double check the software rendering math and lower these // tolerances if possible. - if (format.includes('8unorm')) { + if (format.includes('depth')) { + return 3 / 65536; + } else if (format.includes('8unorm')) { return 7 / 255; } else if (format.includes('2unorm')) { return 9 / 512; @@ -1152,6 +1447,7 @@ export async function putDataInTextureThenDrawAndCheckResultsComparedToSoftwareR const actualTexture = drawTexture(t, texture, samplerDesc, options); const expectedTexelView = softwareRasterize( + t, { descriptor, texels, viewDescriptor }, samplerDesc, [actualTexture.width, actualTexture.height], @@ -1264,7 +1560,7 @@ export function fillTextureWithRandomData(device: GPUDevice, texture: GPUTexture const s_readTextureToRGBA32DeviceToPipeline = new WeakMap< GPUDevice, - Map + Map >(); // MAINTENANCE_TODO: remove cast once textureBindingViewDimension is added to IDL @@ -1296,7 +1592,8 @@ export async function readTextureToTexelViews( s_readTextureToRGBA32DeviceToPipeline.set(device, viewDimensionToPipelineMap); const viewDimension = getEffectiveViewDimension(t, descriptor); - let pipeline = viewDimensionToPipelineMap.get(viewDimension); + const id = `${viewDimension}:${texture.sampleCount}`; + let pipeline = viewDimensionToPipelineMap.get(id); if (!pipeline) { let textureWGSL; let loadWGSL; @@ -1382,7 +1679,7 @@ export async function readTextureToTexelViews( `, }); pipeline = device.createComputePipeline({ layout: 'auto', compute: { module } }); - viewDimensionToPipelineMap.set(viewDimension, pipeline); + viewDimensionToPipelineMap.set(id, pipeline); } const encoder = device.createCommandEncoder(); @@ -1493,6 +1790,14 @@ export async function createTextureWithRandomDataAndGetTexels( } } +function valueIfAllComponentsAreEqual( + c: PerTexelComponent, + componentOrder: TexelComponent[] +) { + const s = new Set(componentOrder.map(component => c[component]!)); + return s.size === 1 ? s.values().next().value : undefined; +} + const kFaceNames = ['+x', '-x', '+y', '-y', '+z', '-z'] as const; /** @@ -1500,14 +1805,19 @@ const kFaceNames = ['+x', '-x', '+y', '-y', '+z', '-z'] as const; * followed by a list of the samples and the weights used for each * component. * - * It works by making an index for every pixel in the texture. Then, - * for each index it generates texture data using TexelView.fromTexelsAsColor - * with a single [1, 1, 1, 1] texel at the texel for the current index. + * It works by making a set of indices for every texel in the texture. + * It splits the set into 2. It picks one set and generates texture data + * using TexelView.fromTexelsAsColor with [1, 1, 1, 1] texels for members + * of the current set. * * In then calls 'run' which renders a single `call`. `run` uses either - * the software renderer or WebGPU. The result ends up being the weights - * used when sampling that pixel. 0 = that texel was not sampled. > 0 = - * it was sampled. + * the software renderer or WebGPU. It then checks the results. If the + * result is zero, all texels in the current had no influence when sampling + * and can be discarded. + * + * If the result is > 0 then, if the set has more than one member, the + * set is split and added to the list to sets to test. If the set only + * had one member then the result is the weight used when sampling that texel. * * This lets you see if the weights from the software renderer match the * weights from WebGPU. @@ -1537,14 +1847,25 @@ const kFaceNames = ['+x', '-x', '+y', '-y', '+z', '-z'] as const; */ async function identifySamplePoints( texture: Texture, - run: (texels: TexelView) => Promise> + run: (texels: TexelView[]) => Promise> ) { const info = texture.descriptor; const isCube = texture.viewDescriptor.dimension === 'cube'; - const textureSize = reifyExtent3D(info.size); - const numTexels = textureSize.width * textureSize.height * textureSize.height; - const texelsPerRow = textureSize.width; - const texelsPerSlice = textureSize.width * textureSize.height; + const mipLevelCount = texture.descriptor.mipLevelCount ?? 1; + const mipLevelSize = range(mipLevelCount, mipLevel => + virtualMipSize(texture.descriptor.dimension ?? '2d', texture.descriptor.size, mipLevel) + ); + const numTexelsPerLevel = mipLevelSize.map(size => size.reduce((s, v) => s * v)); + const numTexelsOfPrecedingLevels = (() => { + let total = 0; + return numTexelsPerLevel.map(v => { + const num = total; + total += v; + return num; + }); + })(); + const numTexels = numTexelsPerLevel.reduce((sum, v) => sum + v); + // This isn't perfect. We already know there was an error. We're just // generating info so it seems okay it's not perfect. This format will // be used to generate weights by drawing with a texture of this format @@ -1583,20 +1904,31 @@ async function identifySamplePoints( } // See if any of the texels in setA were sampled. - const results = await run( - TexelView.fromTexelsAsColors( - format, - (coords: Required): Readonly> => { - const isCandidate = setA.has( - coords.x + coords.y * texelsPerRow + coords.z * texelsPerSlice - ); - const texel: PerTexelComponent = {}; - for (const component of rep.componentOrder) { - texel[component] = isCandidate ? 1 : 0; - } - return texel; - } - ) + const results = convertResultFormatToTexelViewFormat( + await run( + range(mipLevelCount, mipLevel => + TexelView.fromTexelsAsColors( + format, + (coords: Required): Readonly> => { + const size = mipLevelSize[mipLevel]; + const texelsPerSlice = size[0] * size[1]; + const texelsPerRow = size[0]; + const texelId = + numTexelsOfPrecedingLevels[mipLevel] + + coords.x + + coords.y * texelsPerRow + + coords.z * texelsPerSlice; + const isCandidate = setA.has(texelId); + const texel: PerTexelComponent = {}; + for (const component of rep.componentOrder) { + texel[component] = isCandidate ? 1 : 0; + } + return texel; + } + ) + ) + ), + format ); if (rep.componentOrder.some(c => results[c] !== 0)) { // One or more texels of setA were sampled. @@ -1611,6 +1943,31 @@ async function identifySamplePoints( } } + const getMipLevelFromTexelId = (texelId: number) => { + for (let mipLevel = mipLevelCount - 1; mipLevel > 0; --mipLevel) { + if (texelId - numTexelsOfPrecedingLevels[mipLevel] >= 0) { + return mipLevel; + } + } + return 0; + }; + + // separate the sampledTexelWeights by mipLevel, then by layer, within a layer the texelId only includes x and y + const levels: Map>[][] = []; + for (const [texelId, weight] of sampledTexelWeights.entries()) { + const mipLevel = getMipLevelFromTexelId(texelId); + const level = levels[mipLevel] ?? []; + levels[mipLevel] = level; + const size = mipLevelSize[mipLevel]; + const texelsPerSlice = size[0] * size[1]; + const id = texelId - numTexelsOfPrecedingLevels[mipLevel]; + const layer = Math.floor(id / texelsPerSlice); + const layerEntries = level[layer] ?? new Map(); + level[layer] = layerEntries; + const xyId = id - layer * texelsPerSlice; + layerEntries.set(xyId, weight); + } + // ┌───┬───┬───┬───┐ // │ a │ │ │ │ // ├───┼───┼───┼───┤ @@ -1620,66 +1977,92 @@ async function identifySamplePoints( // ├───┼───┼───┼───┤ // │ │ │ │ b │ // └───┴───┴───┴───┘ - const letter = (idx: number) => String.fromCharCode(97 + idx); // 97: 'a' - const orderedTexelIndices: number[] = []; const lines: string[] = []; - for (let z = 0; z < textureSize.depthOrArrayLayers; ++z) { - lines.push(`slice: ${z}${isCube ? ` (${kFaceNames[z]})` : ''}`); - { - let line = ' '; - for (let x = 0; x < textureSize.width; x++) { - line += ` ${x.toString().padEnd(2)}`; - } - lines.push(line); + const letter = (idx: number) => String.fromCodePoint(idx < 30 ? 97 + idx : idx + 9600 - 30); // 97: 'a' + let idCount = 0; + + for (let mipLevel = 0; mipLevel < mipLevelCount; ++mipLevel) { + const level = levels[mipLevel]; + if (!level) { + continue; } - { - let line = ' ┌'; - for (let x = 0; x < textureSize.width; x++) { - line += x === textureSize.width - 1 ? '───┐' : '───┬'; + + const [width, height, depthOrArrayLayers] = mipLevelSize[mipLevel]; + const texelsPerRow = width; + + for (let layer = 0; layer < depthOrArrayLayers; ++layer) { + const layerEntries = level[layer]; + if (!layerEntries) { + continue; } - lines.push(line); - } - for (let y = 0; y < textureSize.height; y++) { + + const orderedTexelIndices: number[] = []; + lines.push(''); + lines.push(`layer: ${layer}${isCube ? ` (${kFaceNames[layer]})` : ''}`); + { - let line = `${y.toString().padEnd(2)}│`; - for (let x = 0; x < textureSize.width; x++) { - const texelIdx = x + y * texelsPerRow + z * texelsPerSlice; - const weight = sampledTexelWeights.get(texelIdx); - if (weight !== undefined) { - line += ` ${letter(orderedTexelIndices.length)} │`; - orderedTexelIndices.push(texelIdx); - } else { - line += ' │'; - } + let line = ' '; + for (let x = 0; x < width; x++) { + line += ` ${x.toString().padEnd(2)}`; } lines.push(line); } - if (y < textureSize.height - 1) { - let line = ' ├'; - for (let x = 0; x < textureSize.width; x++) { - line += x === textureSize.width - 1 ? '───┤' : '───┼'; + { + let line = ' ┌'; + for (let x = 0; x < width; x++) { + line += x === width - 1 ? '───┐' : '───┬'; } lines.push(line); } - } - { - let line = ' └'; - for (let x = 0; x < textureSize.width; x++) { - line += x === textureSize.width - 1 ? '───┘' : '───┴'; + for (let y = 0; y < height; y++) { + { + let line = `${y.toString().padEnd(2)}│`; + for (let x = 0; x < width; x++) { + const texelIdx = x + y * texelsPerRow; + const weight = layerEntries.get(texelIdx); + if (weight !== undefined) { + line += ` ${letter(idCount + orderedTexelIndices.length)} │`; + orderedTexelIndices.push(texelIdx); + } else { + line += ' │'; + } + } + lines.push(line); + } + if (y < height - 1) { + let line = ' ├'; + for (let x = 0; x < width; x++) { + line += x === width - 1 ? '───┤' : '───┼'; + } + lines.push(line); + } } - lines.push(line); + { + let line = ' └'; + for (let x = 0; x < width; x++) { + line += x === width - 1 ? '───┘' : '───┴'; + } + lines.push(line); + } + + const pad2 = (n: number) => n.toString().padStart(2); + const fix5 = (n: number) => n.toFixed(5); + orderedTexelIndices.forEach((texelIdx, i) => { + const weights = layerEntries.get(texelIdx)!; + const y = Math.floor(texelIdx / texelsPerRow); + const x = texelIdx % texelsPerRow; + const singleWeight = valueIfAllComponentsAreEqual(weights, rep.componentOrder); + const w = + singleWeight !== undefined + ? `weight: ${fix5(singleWeight)}` + : `weights: [${rep.componentOrder.map(c => `${c}: ${fix5(weights[c]!)}`).join(', ')}]`; + const coord = `${pad2(x)}, ${pad2(y)}, ${pad2(layer)}`; + lines.push(`${letter(idCount + i)}: mip(${mipLevel}) at: [${coord}], ${w}`); + }); + idCount += orderedTexelIndices.length; } } - const pad2 = (n: number) => n.toString().padStart(2); - orderedTexelIndices.forEach((texelIdx, i) => { - const weights = sampledTexelWeights.get(texelIdx)!; - const z = Math.floor(texelIdx / texelsPerSlice); - const y = Math.floor((texelIdx % texelsPerSlice) / texelsPerRow); - const x = texelIdx % texelsPerRow; - const w = rep.componentOrder.map(c => `${c}: ${weights[c]?.toFixed(5)}`).join(', '); - lines.push(`${letter(i)}: at: [${pad2(x)}, ${pad2(y)}, ${pad2(z)}], weights: [${w}]`); - }); return lines; } @@ -1695,7 +2078,10 @@ function layoutTwoColumns(columnA: string[], columnB: string[]) { return out; } -function getDepthOrArrayLayersForViewDimension(viewDimension?: GPUTextureViewDimension) { +/** + * Returns the number of layers ot test for a given view dimension + */ +export function getDepthOrArrayLayersForViewDimension(viewDimension?: GPUTextureViewDimension) { switch (viewDimension) { case undefined: case '2d': @@ -1728,9 +2114,9 @@ export function chooseTextureSize({ const { blockWidth, blockHeight } = kTextureFormatInfo[format]; const width = align(Math.max(minSize, blockWidth * minBlocks), blockWidth); const height = align(Math.max(minSize, blockHeight * minBlocks), blockHeight); - if (viewDimension === 'cube') { + if (viewDimension === 'cube' || viewDimension === 'cube-array') { const size = lcm(width, height); - return [size, size, 6]; + return [size, size, viewDimension === 'cube-array' ? 24 : 6]; } const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); return [width, height, depthOrArrayLayers]; @@ -1742,15 +2128,6 @@ export type SamplePointMethods = (typeof kSamplePointMethods)[number]; export const kCubeSamplePointMethods = ['cube-edges', 'texel-centre', 'spiral'] as const; export type CubeSamplePointMethods = (typeof kSamplePointMethods)[number]; -/** - * Used to specify a range from [0, num) - * The type is used to determine if values should be integers and if they can be negative. - */ -export type RangeDef = { - num: number; - type: 'f32' | 'i32' | 'u32'; -}; - type TextureBuiltinInputArgs = { descriptor: GPUTextureDescriptor; sampler?: GPUSamplerDescriptor; @@ -1779,6 +2156,7 @@ function generateTextureBuiltinInputsImpl( ): { coords: T; mipLevel: number; sampleIndex?: number; arrayIndex?: number; offset?: T }[] { const { method, descriptor } = args; const dimension = descriptor.dimension ?? '2d'; + const mipLevelCount = descriptor.mipLevelCount ?? 1; const size = virtualMipSize(dimension, descriptor.size, 0); const coords: T[] = []; switch (method) { @@ -1829,8 +2207,11 @@ function generateTextureBuiltinInputsImpl( const kSubdivisionsPerTexel = 4; const nearest = !args.sampler || args.sampler.minFilter === 'nearest'; return coords.map((c, i) => { - const mipLevel = args.mipLevel ? makeRangeValue(args.mipLevel, i) : 0; - const mipSize = virtualMipSize(dimension, size, mipLevel); + const mipLevel = args.mipLevel + ? quantizeMipLevel(makeRangeValue(args.mipLevel, i), args.sampler?.mipmapFilter ?? 'nearest') + : 0; + const clampedMipLevel = clamp(mipLevel, { min: 0, max: mipLevelCount - 1 }); + const mipSize = virtualMipSize(dimension, size, clampedMipLevel); const q = mipSize.map(v => v * kSubdivisionsPerTexel); const coords = c.map((v, i) => { @@ -1855,6 +2236,20 @@ function generateTextureBuiltinInputsImpl( }); } +const kMipEpsilon = 0.02; +function quantizeMipLevel(mipLevel: number, mipmapFilter: GPUFilterMode) { + if (mipmapFilter === 'linear') { + return mipLevel; + } + const intMip = Math.floor(mipLevel); + const fractionalMip = mipLevel - intMip; + if (fractionalMip < 0.5 - kMipEpsilon || fractionalMip > 0.5 + kMipEpsilon) { + return mipLevel; + } else { + return intMip + 0.5 + (fractionalMip < 0.5 ? -kMipEpsilon : +kMipEpsilon); + } +} + // Removes the first element from an array of types type FilterFirstElement = T extends [unknown, ...infer R] ? R : []; @@ -1916,9 +2311,8 @@ function normalize(v: vec3): vec3 { /** * Converts a cube map coordinate to a uv coordinate (0 to 1) and layer (0.5/6.0 to 5.5/6.0). - * Also returns the length of the original coordinate. */ -function convertCubeCoordToNormalized3DTextureCoord(v: vec3): vec3 { +export function convertCubeCoordToNormalized3DTextureCoord(v: vec3): vec3 { let uvw; let layer; // normalize the coord. @@ -1947,7 +2341,7 @@ function convertCubeCoordToNormalized3DTextureCoord(v: vec3): vec3 { /** * Convert a 3d texcoord into a cube map coordinate. */ -function convertNormalized3DTexCoordToCubeCoord(uvLayer: vec3) { +export function convertNormalized3DTexCoordToCubeCoord(uvLayer: vec3) { const [u, v, faceLayer] = uvLayer; return normalize(transformMat3([u, v, 1], kFaceUVMatrices[Math.min(5, faceLayer * 6) | 0])); } @@ -2022,7 +2416,7 @@ const kFaceToFaceRemap: { to: number; u: FaceCoordConversion; v: FaceCoordConver [ /* -u */ { to: 1, u: 'v', v: '1+u' }, /* +u */ { to: 0, u: 't-v-1', v: 'u-t' }, - /* -v */ { to: 5, u: 't-u-1', v: 't-v-1' }, + /* -v */ { to: 5, u: 't-u-1', v: '-v-1' }, /* +v */ { to: 4, u: 'u', v: 'v-t' }, ], // 3 @@ -2079,9 +2473,9 @@ function applyFaceWrap(textureSize: number, faceCoord: vec3): vec3 { function wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize: number, faceCoord: vec3) { // If we're off both edges we need to wrap twice, once for each edge. - faceCoord = applyFaceWrap(textureSize, faceCoord); - faceCoord = applyFaceWrap(textureSize, faceCoord); - return faceCoord; + const faceCoord1 = applyFaceWrap(textureSize, faceCoord); + const faceCoord2 = applyFaceWrap(textureSize, faceCoord1); + return faceCoord2; } function applyAddressModesToCoords( @@ -2130,6 +2524,7 @@ export function generateSamplePointsCube( offset?: undefined; }[] { const { method, descriptor } = args; + const mipLevelCount = descriptor.mipLevelCount ?? 1; const size = virtualMipSize('2d', descriptor.size, 0); const textureWidth = size[0]; const coords: vec3[] = []; @@ -2219,7 +2614,8 @@ export function generateSamplePointsCube( const nearest = !args.sampler || args.sampler.minFilter === 'nearest'; return coords.map((c, i) => { const mipLevel = args.mipLevel ? makeRangeValue(args.mipLevel, i) : 0; - const mipSize = virtualMipSize('2d', size, mipLevel); + const clampedMipLevel = clamp(mipLevel, { min: 0, max: mipLevelCount - 1 }); + const mipSize = virtualMipSize('2d', size, clampedMipLevel); const q = [ mipSize[0] * kSubdivisionsPerTexel, mipSize[0] * kSubdivisionsPerTexel, @@ -2230,9 +2626,9 @@ export function generateSamplePointsCube( // If this is a corner, move to in so it's not // (see comment "Issues with corners of cubemaps") - const ndx = getUnusedCubeCornerSampleIndex(textureWidth, uvw); + const ndx = getUnusedCubeCornerSampleIndex(mipSize[0], uvw); if (ndx >= 0) { - const halfTexel = 0.5 / textureWidth; + const halfTexel = 0.5 / mipSize[0]; uvw[0] = clamp(uvw[0], { min: halfTexel, max: 1 - halfTexel }); } @@ -2357,12 +2753,8 @@ function buildBinnedCalls(calls: TextureCall[]) { 'texture calls are not binned correctly' ); if (value !== undefined && name !== 'offset') { - const bitcastToU32 = (value: number) => { - if (calls[0].coordType === 'f') { - return float32ToUint32(value); - } - return value; - }; + const type = getCallArgType(call, name); + const bitcastToU32 = kBitCastFunctions[type]; if (value instanceof Array) { for (const c of value) { data.push(bitcastToU32(c)); @@ -2521,11 +2913,12 @@ ${body} const pipelines = s_deviceToPipelines.get(t.device) ?? new Map(); s_deviceToPipelines.set(t.device, pipelines); - let pipeline = pipelines.get(code); + const id = `${renderTarget.format}:${code}`; + let pipeline = pipelines.get(id); if (!pipeline) { const shaderModule = t.device.createShaderModule({ code }); - pipeline = t.device.createRenderPipeline({ + pipeline = await t.device.createRenderPipelineAsync({ layout: 'auto', vertex: { module: shaderModule }, fragment: { @@ -2535,7 +2928,7 @@ ${body} primitive: { topology: 'triangle-strip' }, }); - pipelines.set(code, pipeline); + pipelines.set(id, pipeline); } const gpuSampler = sampler ? t.device.createSampler(sampler) : undefined; diff --git a/src/webgpu/util/texture.ts b/src/webgpu/util/texture.ts index f1423ce2fd63..badce71baa34 100644 --- a/src/webgpu/util/texture.ts +++ b/src/webgpu/util/texture.ts @@ -294,6 +294,11 @@ function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { `; } +const s_copyBufferToTextureViaRenderPipelines = new WeakMap< + GPUDevice, + Map +>(); + function copyBufferToTextureViaRender( t: GPUTest, encoder: GPUCommandEncoder, @@ -311,33 +316,41 @@ function copyBufferToTextureViaRender( const { device } = t; const code = getCopyBufferToTextureViaRenderCode(format); - const module = device.createShaderModule({ code }); - const pipeline = device.createRenderPipeline({ - layout: 'auto', - vertex: { module }, - ...(useFragDepth - ? { - fragment: { - module, - targets: [], - }, - depthStencil: { - depthWriteEnabled: true, - depthCompare: 'always', - format, - }, - } - : { - fragment: { - module, - targets: [{ format }], - }, - }), - primitive: { - topology: 'triangle-strip', - }, - ...(sampleCount > 1 && { multisample: { count: sampleCount } }), - }); + const id = JSON.stringify({ format, useFragDepth, sampleCount, code }); + const pipelines = + s_copyBufferToTextureViaRenderPipelines.get(device) ?? new Map(); + s_copyBufferToTextureViaRenderPipelines.set(device, pipelines); + let pipeline = pipelines.get(id); + if (!pipeline) { + const module = device.createShaderModule({ code }); + pipeline = device.createRenderPipeline({ + layout: 'auto', + vertex: { module }, + ...(useFragDepth + ? { + fragment: { + module, + targets: [], + }, + depthStencil: { + depthWriteEnabled: true, + depthCompare: 'always', + format, + }, + } + : { + fragment: { + module, + targets: [{ format }], + }, + }), + primitive: { + topology: 'triangle-strip', + }, + ...(sampleCount > 1 && { multisample: { count: sampleCount } }), + }); + pipelines.set(id, pipeline); + } const info = kTextureFormatInfo[format]; const uniforms = new Uint32Array([ @@ -371,7 +384,7 @@ function copyBufferToTextureViaRender( mipLevelCount, arrayLayerCount, }), - depthClearValue: 0.5, + depthClearValue: 0, depthLoadOp: 'clear', depthStoreOp: 'store', },