From 2a86881ebbf57050b5a092f48ea907926a291ab9 Mon Sep 17 00:00:00 2001 From: Gregg Tavares Date: Mon, 7 Oct 2024 18:06:47 +0900 Subject: [PATCH] WGSL texture builtins all stages Test texture builtins on all stages. Previously only the fragment stage was tested. Note: Some of these are expected to fail on Intel Mac because in compute shaders, Intel Mac doesn't do bilinear interpolation between mip levels. At least not if not using argument buffers. --- .../call/builtin/textureGather.spec.ts | 132 +++- .../call/builtin/textureGatherCompare.spec.ts | 76 ++- .../call/builtin/textureLoad.spec.ts | 180 ++++-- .../textureSampleBaseClampToEdge.spec.ts | 17 +- .../builtin/textureSampleCompareLevel.spec.ts | 76 ++- .../call/builtin/textureSampleGrad.spec.ts | 65 +- .../call/builtin/textureSampleLevel.spec.ts | 117 +++- .../expression/call/builtin/texture_utils.ts | 574 ++++++++++++++---- 8 files changed, 982 insertions(+), 255 deletions(-) diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts index fddb65e1420f..88f51dfdfe0c 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts @@ -32,6 +32,7 @@ import { kDepthStencilFormats, kEncodableTextureFormats, } from '../../../../../format_info.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { appendComponentTypeForFormatToTextureType, @@ -83,6 +84,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isFillable(t.format)) .combine('minFilter', ['nearest', 'linear'] as const) @@ -98,7 +100,8 @@ Parameters: skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); }) .fn(async t => { - const { format, C, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, C, samplePoints, stage, 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 }); @@ -137,14 +140,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -169,6 +181,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isFillable(t.format)) .combine('minFilter', ['nearest', 'linear'] as const) @@ -182,7 +195,7 @@ Parameters: skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); }) .fn(async t => { - const { format, C, samplePoints, addressMode, minFilter } = t.params; + const { format, C, stage, samplePoints, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube'; const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -225,14 +238,23 @@ Parameters: dimension: viewDimension, }; const textureType = appendComponentTypeForFormatToTextureType('texture_cube', format); - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -266,6 +288,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isFillable(t.format)) .combine('minFilter', ['nearest', 'linear'] as const) @@ -282,7 +305,8 @@ Parameters: skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); }) .fn(async t => { - const { format, samplePoints, C, A, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, samplePoints, C, 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 }); @@ -326,14 +350,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_2d_array', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -361,6 +394,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isFillable(t.format)) .combine('minFilter', ['nearest', 'linear'] as const) @@ -376,7 +410,7 @@ Parameters: skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); }) .fn(async t => { - const { format, C, A, samplePoints, addressMode, minFilter } = t.params; + const { format, C, A, stage, samplePoints, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -421,14 +455,23 @@ Parameters: dimension: viewDimension, }; const textureType = appendComponentTypeForFormatToTextureType('texture_cube_array', format); - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -454,6 +497,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -467,7 +511,7 @@ Parameters: .combine('offset', [false, true] as const) ) .fn(async t => { - const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -503,14 +547,23 @@ Parameters: }); const textureType = 'texture_depth_2d'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -529,6 +582,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -540,7 +594,7 @@ Parameters: .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) ) .fn(async t => { - const { format, samplePoints, addressMode, minFilter } = t.params; + const { format, stage, samplePoints, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube'; const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -580,14 +634,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -616,6 +679,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -634,7 +698,8 @@ Parameters: skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); }) .fn(async t => { - const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -675,14 +740,23 @@ Parameters: }); const textureType = 'texture_depth_2d_array'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -704,6 +778,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -719,7 +794,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported('cube-array'); }) .fn(async t => { - const { format, A, samplePoints, addressMode, minFilter } = t.params; + const { format, A, stage, samplePoints, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -761,14 +836,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube_array'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts index cad0e85c6625..742badde94af 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts @@ -23,6 +23,7 @@ import { isEncodableTextureFormat, kDepthStencilFormats, } from '../../../../../format_info.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { checkCallResults, @@ -67,6 +68,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -85,8 +87,17 @@ Parameters: t.skipIfTextureFormatNotSupported(t.params.format); }) .fn(async t => { - const { format, samplePoints, A, addressModeU, addressModeV, minFilter, compare, offset } = - t.params; + const { + format, + stage, + samplePoints, + A, + addressModeU, + addressModeV, + minFilter, + compare, + offset, + } = t.params; const viewDimension = '2d-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension }); @@ -130,14 +141,23 @@ Parameters: }); const textureType = 'texture_depth_2d_array'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -160,6 +180,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -176,7 +197,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported('cube-array'); }) .fn(async t => { - const { format, A, samplePoints, addressMode, minFilter, compare } = t.params; + const { format, A, stage, samplePoints, addressMode, minFilter, compare } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -222,14 +243,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube_array'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -256,6 +286,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -270,7 +301,7 @@ Parameters: .combine('offset', [false, true] as const) ) .fn(async t => { - const { format, C, samplePoints, addressMode, compare, minFilter, offset } = t.params; + const { format, C, stage, samplePoints, addressMode, compare, minFilter, offset } = t.params; const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); const descriptor: GPUTextureDescriptor = { @@ -309,14 +340,23 @@ Parameters: }); const textureType = 'texture_depth_2d'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -336,6 +376,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -348,7 +389,7 @@ Parameters: .combine('compare', kCompareFunctions) ) .fn(async t => { - const { format, samplePoints, addressMode, minFilter, compare } = t.params; + const { format, stage, samplePoints, addressMode, minFilter, compare } = t.params; const viewDimension: GPUTextureViewDimension = 'cube'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -391,14 +432,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts index 2f42fffeb768..9bb4f74b035b 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts @@ -38,6 +38,7 @@ import { import { GPUTest } from '../../../../../gpu_test.js'; import { maxMipLevelCount, virtualMipSize } from '../../../../../util/texture/base.js'; import { TexelFormats } from '../../../../types.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { TextureCall, @@ -90,6 +91,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => textureDimensionAndFormatCompatible('1d', t.format)) // 1d textures can't have a height !== 1 @@ -105,7 +107,7 @@ Parameters: t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, C, L, samplePoints } = t.params; + const { format, stage, C, L, samplePoints } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const [width] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -137,14 +139,23 @@ Parameters: const textureType = appendComponentTypeForFormatToTextureType('texture_1d', texture.format); const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -166,6 +177,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => !isCompressedFloatTextureFormat(t.format)) .beginSubcases() @@ -179,7 +191,7 @@ Parameters: t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, samplePoints, C, L } = t.params; + const { format, stage, samplePoints, C, L } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -209,14 +221,23 @@ Parameters: const textureType = appendComponentTypeForFormatToTextureType('texture_2d', texture.format); const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -237,6 +258,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => textureDimensionAndFormatCompatible('3d', t.format)) .beginSubcases() @@ -250,7 +272,7 @@ Parameters: t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { format, samplePoints, C, L } = t.params; + const { format, stage, samplePoints, C, L } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); @@ -282,14 +304,23 @@ Parameters: const textureType = appendComponentTypeForFormatToTextureType('texture_3d', texture.format); const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -312,6 +343,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('texture_type', [ 'texture_multisampled_2d', 'texture_depth_multisampled_2d', @@ -336,7 +368,7 @@ Parameters: t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { texture_type, format, samplePoints, C, S } = t.params; + const { texture_type, format, stage, samplePoints, C, S } = t.params; const sampleCount = 4; const descriptor: GPUTextureDescriptor = { @@ -368,14 +400,23 @@ Parameters: const textureType = appendComponentTypeForFormatToTextureType(texture_type, texture.format); const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -396,6 +437,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -410,7 +452,7 @@ Parameters: t.skipIfTextureLoadNotSupportedForTextureType('texture_depth_2d'); }) .fn(async t => { - const { format, samplePoints, C, L } = t.params; + const { format, stage, samplePoints, C, L } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -440,14 +482,23 @@ Parameters: const textureType = 'texture_depth_2d'; const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -465,8 +516,10 @@ Parameters: * coords: The 0-based texel coordinate. ` ) - .paramsSubcasesOnly(u => + .params(u => u + .combine('stage', kShaderStages) + .beginSubcases() .combine('samplePoints', kSamplePointMethods) .combine('C', ['i32', 'u32'] as const) .combine('L', ['i32', 'u32'] as const) @@ -475,7 +528,7 @@ Parameters: t.skipIf(typeof VideoFrame === 'undefined', 'VideoFrames are not supported') ) .fn(async t => { - const { samplePoints, C, L } = t.params; + const { stage, samplePoints, C, L } = t.params; const size = [8, 8, 1]; @@ -505,14 +558,23 @@ Parameters: const textureType = 'texture_external'; const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); videoFrame.close(); @@ -536,6 +598,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) // MAINTENANCE_TODO: Update createTextureFromTexelViews to support stencil8 and remove this filter. .filter(t => t.format !== 'stencil8' && !isCompressedFloatTextureFormat(t.format)) @@ -559,7 +622,7 @@ Parameters: t.selectDeviceForTextureFormatOrSkipTestCase(t.params.format); }) .fn(async t => { - const { texture_type, format, samplePoints, C, A, L } = t.params; + const { texture_type, format, stage, samplePoints, C, A, L } = t.params; // We want at least 4 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); @@ -595,14 +658,23 @@ Parameters: const textureType = appendComponentTypeForFormatToTextureType(texture_type, texture.format); const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -622,6 +694,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -636,7 +709,7 @@ Parameters: } }) .fn(async t => { - const { format, samplePoints, C } = t.params; + const { format, stage, samplePoints, C } = t.params; // We want at least 3 blocks or something wide enough for 3 mip levels. const [width] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); @@ -663,14 +736,23 @@ Parameters: const textureType = `texture_storage_1d<${format}, read>`; const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -690,6 +772,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -704,7 +787,7 @@ Parameters: } }) .fn(async t => { - const { format, samplePoints, C } = t.params; + const { format, stage, samplePoints, C } = t.params; // We want at least 3 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 3, format }); @@ -729,14 +812,23 @@ Parameters: const textureType = `texture_storage_2d<${format}, read>`; const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -758,6 +850,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -773,7 +866,7 @@ Parameters: } }) .fn(async t => { - const { format, samplePoints, C, A } = t.params; + const { format, stage, samplePoints, C, A } = t.params; // We want at least 3 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); @@ -803,14 +896,23 @@ Parameters: dimension: '2d-array', }; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -830,6 +932,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combineWithParams([...TexelFormats, { format: 'bgra8unorm' }] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -844,7 +947,7 @@ Parameters: } }) .fn(async t => { - const { format, samplePoints, C } = t.params; + const { format, stage, samplePoints, C } = t.params; // We want at least 3 blocks or something wide enough for 3 mip levels. const size = chooseTextureSize({ minSize: 8, minBlocks: 4, format, viewDimension: '3d' }); @@ -870,14 +973,23 @@ Parameters: const textureType = `texture_storage_3d<${format}, read>`; const viewDescriptor = {}; const sampler = undefined; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts index b670c44035f9..e3bf2c01b136 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleBaseClampToEdge.spec.ts @@ -5,6 +5,7 @@ Execution tests for textureSampleBaseClampToEdge import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { GPUTest } from '../../../../../gpu_test.js'; import { TexelView } from '../../../../../util/texture/texel_view.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { checkCallResults, @@ -54,6 +55,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('textureType', ['texture_2d', 'texture_external'] as const) .beginSubcases() .combine('samplePoints', kSamplePointMethods) @@ -68,7 +70,7 @@ Parameters: ) ) .fn(async t => { - const { textureType, samplePoints, addressModeU, addressModeV, minFilter } = t.params; + const { textureType, stage, samplePoints, addressModeU, addressModeV, minFilter } = t.params; const descriptor: GPUTextureDescriptor = { format: 'rgba8unorm', @@ -105,14 +107,23 @@ Parameters: }; }); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); } finally { diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts index 75077b5afa74..1062a76a4c7f 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleCompareLevel.spec.ts @@ -18,6 +18,7 @@ import { isEncodableTextureFormat, kDepthStencilFormats, } from '../../../../../format_info.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { checkCallResults, @@ -59,6 +60,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -73,7 +75,7 @@ Parameters: .combine('offset', [false, true] as const) ) .fn(async t => { - const { format, samplePoints, addressModeU, addressModeV, minFilter, compare, offset } = + const { format, stage, samplePoints, addressModeU, addressModeV, minFilter, compare, offset } = t.params; const size = chooseTextureSize({ minSize: 16, minBlocks: 4, format }); @@ -117,14 +119,23 @@ Parameters: }); const textureType = 'texture_depth_2d'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -144,6 +155,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -156,7 +168,7 @@ Parameters: .combine('compare', kCompareFunctions) ) .fn(async t => { - const { format, samplePoints, addressMode, minFilter, compare } = t.params; + const { format, stage, samplePoints, addressMode, minFilter, compare } = t.params; const viewDimension: GPUTextureViewDimension = 'cube'; const size = chooseTextureSize({ minSize: 16, minBlocks: 2, format, viewDimension }); @@ -203,14 +215,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -240,6 +261,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -258,8 +280,17 @@ Parameters: t.skipIfTextureFormatNotSupported(t.params.format); }) .fn(async t => { - const { format, samplePoints, A, addressModeU, addressModeV, minFilter, compare, offset } = - t.params; + const { + format, + stage, + samplePoints, + A, + addressModeU, + addressModeV, + minFilter, + compare, + offset, + } = t.params; const viewDimension = '2d-array'; const size = chooseTextureSize({ minSize: 16, minBlocks: 4, format, viewDimension }); @@ -306,14 +337,23 @@ Parameters: }); const textureType = 'texture_depth_2d_array'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -336,6 +376,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -352,7 +393,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported('cube-array'); }) .fn(async t => { - const { format, A, samplePoints, addressMode, minFilter, compare } = t.params; + const { format, A, stage, samplePoints, addressMode, minFilter, compare } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); @@ -400,14 +441,23 @@ Parameters: dimension: viewDimension, }; const textureType = 'texture_depth_cube_array'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); diff --git a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts index bbff69d32969..774fe0f16cc1 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleGrad.spec.ts @@ -7,6 +7,7 @@ Samples a texture using explicit gradients. import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { kCompressedTextureFormats, kEncodableTextureFormats } from '../../../../../format_info.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { appendComponentTypeForFormatToTextureType, @@ -57,6 +58,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -70,7 +72,7 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -108,14 +110,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -144,6 +155,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('viewDimension', ['3d', 'cube'] as const) @@ -165,6 +177,7 @@ Parameters: const { format, viewDimension, + stage, samplePoints, addressModeU, addressModeV, @@ -240,14 +253,23 @@ Parameters: dimension: viewDimension, }; const textureType = getTextureTypeForTextureViewDimension(viewDimension)!; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -278,6 +300,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -292,7 +315,8 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -335,14 +359,23 @@ Parameters: }); const textureType = 'texture_2d_array'; const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -372,6 +405,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -385,7 +419,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported('cube-array'); }) .fn(async t => { - const { format, samplePoints, A, addressMode, minFilter } = t.params; + const { format, stage, samplePoints, A, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ @@ -433,14 +467,23 @@ Parameters: dimension: viewDimension, }; const textureType = getTextureTypeForTextureViewDimension(viewDimension); - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); 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 5c6e99eb9665..14f6b936a429 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts @@ -1,9 +1,6 @@ export const description = ` 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. `; @@ -16,6 +13,7 @@ import { kDepthStencilFormats, kEncodableTextureFormats, } from '../../../../../format_info.js'; +import { kShaderStages } from '../../../../validation/decl/util.js'; import { appendComponentTypeForFormatToTextureType, @@ -69,6 +67,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -82,7 +81,7 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -120,14 +119,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_2d', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -161,6 +169,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -175,7 +184,8 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, A, addressModeU, addressModeV, minFilter, offset } = t.params; + const { format, stage, 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 }); @@ -218,14 +228,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_2d_array', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -257,6 +276,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .combine('viewDimension', ['3d', 'cube'] as const) @@ -273,7 +293,7 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, viewDimension, samplePoints, addressMode, minFilter, offset } = t.params; + const { format, viewDimension, stage, samplePoints, addressMode, minFilter, offset } = t.params; const [width, height] = chooseTextureSize({ minSize: 32, minBlocks: 2, format, viewDimension }); const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); @@ -327,14 +347,23 @@ Parameters: dimension: viewDimension, }; const textureType = getTextureTypeForTextureViewDimension(viewDimension); - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -369,6 +398,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kTestableColorFormats) .filter(t => isPotentiallyFilterableAndFillable(t.format)) .beginSubcases() @@ -382,7 +412,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported('cube-array'); }) .fn(async t => { - const { format, samplePoints, A, addressMode, minFilter } = t.params; + const { format, stage, samplePoints, A, addressMode, minFilter } = t.params; const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ @@ -429,14 +459,23 @@ Parameters: dimension: viewDimension, }; const textureType = getTextureTypeForTextureViewDimension(viewDimension); - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -469,6 +508,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -485,7 +525,7 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, addressMode, minFilter, L, offset } = t.params; + const { format, stage, 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 }); @@ -523,14 +563,23 @@ Parameters: }); const textureType = appendComponentTypeForFormatToTextureType('texture_depth_2d', format); const viewDescriptor = {}; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -565,6 +614,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -582,7 +632,7 @@ Parameters: skipIfTextureFormatNotSupportedNotAvailableOrNotFilterable(t, t.params.format) ) .fn(async t => { - const { format, samplePoints, addressMode, minFilter, A, L, offset } = t.params; + const { format, stage, 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 }); @@ -624,14 +674,23 @@ Parameters: }); 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 results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, { texels, descriptor, viewDescriptor }, textureType, sampler, calls, - results + results, + stage ); t.expectOK(res); }); @@ -665,6 +724,7 @@ Parameters: ) .params(u => u + .combine('stage', kShaderStages) .combine('format', kDepthStencilFormats) // filter out stencil only formats .filter(t => isDepthTextureFormat(t.format)) @@ -686,7 +746,7 @@ Parameters: t.skipIfTextureViewDimensionNotSupported(t.params.viewDimension); }) .fn(async t => { - const { format, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params; + const { format, stage, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params; const size = chooseTextureSize({ minSize: 32, @@ -734,7 +794,15 @@ Parameters: }; const textureType = viewDimension === 'cube' ? 'texture_depth_cube' : 'texture_depth_cube_array'; - const results = await doTextureCalls(t, texture, viewDescriptor, textureType, sampler, calls); + const results = await doTextureCalls( + t, + texture, + viewDescriptor, + textureType, + sampler, + calls, + stage + ); const res = await checkCallResults( t, @@ -742,7 +810,8 @@ Parameters: textureType, sampler, calls, - results + results, + stage ); 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 2f64739c86c4..b706a7add377 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -40,6 +40,7 @@ import { import { PerPixelAtLevel, TexelView } from '../../../../../util/texture/texel_view.js'; import { createTextureFromTexelViews } from '../../../../../util/texture.js'; import { reifyExtent3D } from '../../../../../util/unions.js'; +import { ShaderStage } from '../../../../validation/decl/util.js'; export const kSampleTypeInfo = { f32: { @@ -164,22 +165,66 @@ type MipWeights = { }; type MipWeightType = keyof MipWeights; +function makeGraph(width: number, height: number) { + const data = new Uint8Array(width * height); + + return { + plot(norm: number, x: number, c: number) { + const y = clamp(Math.round(norm * height), { min: 0, max: height - 1 }); + const offset = (height - y - 1) * width + x; + data[offset] = c; + }, + plotValues(values: Iterable, c: number) { + let i = 0; + for (const v of values) { + this.plot(v, i, c); + ++i; + } + }, + toString(conversion = ['.', 'e', 'A']) { + const lines = []; + for (let y = 0; y < height; ++y) { + const offset = y * width; + lines.push([...data.subarray(offset, offset + width)].map(v => conversion[v]).join('')); + } + return lines.join('\n'); + }, + }; +} + +function* linear0to1OverN(n: number) { + for (let i = 0; i <= n; ++i) { + yield i / n; + } +} + +function graphWeights(height: number, weights: number[]) { + const graph = makeGraph(weights.length, height); + graph.plotValues(linear0to1OverN(weights.length - 1), 1); + graph.plotValues(weights, 2); + return graph.toString(); +} + /** * Validates the weights go from 0 to 1 in increasing order. */ -function validateWeights(weights: number[]) { - const showWeights = () => weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n'); +function validateWeights(stage: string, weights: number[]) { + const showWeights = () => ` +${weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n')} + +e = expected +A = actual +${graphWeights(32, weights)} +`; // Validate the weights - assert(weights[0] === 0, `weight 0 expected 0 but was ${weights[0]}\n${showWeights()}`); assert( - weights[kMipGradientSteps] === 1, - `top weight expected 1 but was ${weights[kMipGradientSteps]}\n${showWeights()}` + weights[0] === 0, + `stage: ${stage}, weight 0 expected 0 but was ${weights[0]}\n${showWeights()}` ); - assert( - new Set(weights).size >= ((weights.length * 0.66) | 0), - `expected more unique weights\n${showWeights()}` + weights[kMipGradientSteps] === 1, + `stage: ${stage}, top weight expected 1 but was ${weights[kMipGradientSteps]}\n${showWeights()}` ); // Note: for 16 steps, these are the AMD weights @@ -207,12 +252,68 @@ function validateWeights(weights: number[]) { // // notice step 1 is 0 and step 15 is 1. // so we only check the 1 through 14. + // + // Note: these 2 changes are effectively here to catch Intel Mac + // issues and require implementations to work around them. + // + // Ideally the weights should form a straight line + // + // +----------------+ + // | **| + // | ** | + // | ** | + // | ** | + // | ** | + // | ** | + // | ** | + // |** | + // +----------------+ + // + // AMD Mac goes like this: Not great but we allow it + // + // +----------------+ + // | ***| + // | ** | + // | * | + // | ** | + // | ** | + // | * | + // | ** | + // |*** | + // +----------------+ + // + // Intel Mac goes like this: Unacceptable + // + // +----------------+ + // | *******| + // | * | + // | * | + // | * | + // | * | + // | * | + // | * | + // |******* | + // +----------------+ + // + const dx = 1 / kMipGradientSteps; for (let i = 0; i < kMipGradientSteps; ++i) { + const dy = weights[i + 1] - weights[i]; + // dy / dx because dy might be 0 + const slope = dy / dx; + assert( + slope >= 0, + `stage: ${stage}, weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}` + ); assert( - weights[i] <= weights[i + 1], - `weight[${i}] was not <= weight[${i + 1}]\n${showWeights()}` + slope <= 2, + `stage: ${stage}, slope from weight[${i}] to weight[${i + 1}] is > 2.\n${showWeights()}` ); } + + assert( + new Set(weights).size >= ((weights.length * 0.66) | 0), + `stage: ${stage}, expected more unique weights\n${showWeights()}` + ); } /** @@ -328,43 +429,78 @@ async function queryMipGradientValuesForDevice(t: GPUTest) { struct VSOutput { @builtin(position) pos: vec4f, @location(0) @interpolate(flat, either) ndx: u32, + @location(1) @interpolate(flat, either) result: vec4f, }; - @vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput { + fn getMixLevels(wNdx: u32) -> vec4f { + let mipLevel = f32(wNdx) / ${kMipGradientSteps}; + let size = textureDimensions(tex); + let g = mix(1.0, 2.0, mipLevel) / f32(size.x); + let ddx = vec2f(g, 0); + return vec4f( + textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r, + textureSampleGrad(tex, smp, vec2f(0.5), ddx, vec2f(0)).r, + 0, + 0); + } + + fn recordMixLevels(wNdx: u32, r: vec4f) { + let ndx = wNdx * ${kNumWeightTypes}; + for (var i: u32 = 0; i < ${kNumWeightTypes}; i++) { + result[ndx + i] = r[i]; + } + } + + fn getPosition(vNdx: u32) -> vec4f { let pos = array( vec2f(-1, 3), vec2f( 3, -1), vec2f(-1, -1), ); let p = pos[vNdx]; - return VSOutput(vec4f(p, 0, 1), iNdx); + return vec4f(p, 0, 1); } - @fragment fn fs(v: VSOutput) -> @location(0) vec4f { - let mipLevel = f32(v.ndx) / ${kMipGradientSteps}; - let size = textureDimensions(tex); - let g = mix(1.0, 2.0, mipLevel) / f32(size.x); - let ddx = vec2f(g, 0); + @vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput { + return VSOutput(getPosition(vNdx), iNdx, vec4f(0)); + } + + @fragment fn fsRecord(v: VSOutput) -> @location(0) vec4f { + recordMixLevels(v.ndx, getMixLevels(v.ndx)); + return vec4f(0); + } + + @compute @workgroup_size(1) fn csRecord(@builtin(global_invocation_id) id: vec3u) { + recordMixLevels(id.x, getMixLevels(id.x)); + } + + @vertex fn vsRecord(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput { + return VSOutput(getPosition(vNdx), iNdx, getMixLevels(iNdx)); - let ndx = v.ndx * ${kNumWeightTypes}; - result[ndx + 0] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r; - result[ndx + 1] = textureSampleGrad(tex, smp, vec2f(0.5), ddx, vec2f(0)).r; + } + @fragment fn fsSaveVs(v: VSOutput) -> @location(0) vec4f { + recordMixLevels(v.ndx, v.result); return vec4f(0); } `, }); - const pipeline = device.createRenderPipeline({ + const vertexPipeline = device.createRenderPipeline({ layout: 'auto', - vertex: { module }, - fragment: { module, targets: [{ format: 'rgba8unorm' }] }, + vertex: { module, entryPoint: 'vsRecord' }, + fragment: { module, entryPoint: 'fsSaveVs', targets: [{ format: 'rgba8unorm' }] }, }); - const target = t.createTextureTracked({ - size: [1, 1], - format: 'rgba8unorm', - usage: GPUTextureUsage.RENDER_ATTACHMENT, + const fragmentPipeline = device.createRenderPipeline({ + layout: 'auto', + vertex: { module, entryPoint: 'vs' }, + fragment: { module, entryPoint: 'fsRecord', targets: [{ format: 'rgba8unorm' }] }, + }); + + const computePipeline = device.createComputePipeline({ + layout: 'auto', + compute: { module }, }); const texture = t.createTextureTracked({ @@ -387,59 +523,131 @@ async function queryMipGradientValuesForDevice(t: GPUTest) { mipmapFilter: 'linear', }); + const target = t.createTextureTracked({ + size: [1, 1], + format: 'rgba8unorm', + usage: GPUTextureUsage.RENDER_ATTACHMENT, + }); + const storageBuffer = t.createBufferTracked({ size: 4 * (kMipGradientSteps + 1) * kNumWeightTypes, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); - const resultBuffer = t.createBufferTracked({ - size: storageBuffer.size, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, - }); + type PassFn = ( + encoder: GPUCommandEncoder, + bindGroup: GPUBindGroup, + resultBuffer: GPUBuffer + ) => void; + + const getMixWeightForStage = ( + encoder: GPUCommandEncoder, + pipeline: GPUComputePipeline | GPURenderPipeline, + passFn: PassFn + ) => { + 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 bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: texture.createView() }, + { binding: 1, resource: sampler }, + { binding: 2, resource: { buffer: storageBuffer } }, + ], + }); + + passFn(encoder, bindGroup, resultBuffer); + encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size); + return resultBuffer; + }; const encoder = device.createCommandEncoder(); - const pass = encoder.beginRenderPass({ - colorAttachments: [ - { - view: target.createView(), - loadOp: 'clear', - storeOp: 'store', - }, - ], - }); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bindGroup); - pass.draw(3, kMipGradientSteps + 1); - pass.end(); - encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size); + const stageBuffers: Record = { + compute: getMixWeightForStage( + encoder, + computePipeline, + (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => { + const pass = encoder.beginComputePass(); + pass.setPipeline(computePipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(kMipGradientSteps + 1); + pass.end(); + } + ), + fragment: getMixWeightForStage( + encoder, + fragmentPipeline, + (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => { + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: target.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(fragmentPipeline); + pass.setBindGroup(0, bindGroup); + pass.draw(3, kMipGradientSteps + 1); + pass.end(); + } + ), + vertex: getMixWeightForStage( + encoder, + vertexPipeline, + (encoder: GPUCommandEncoder, bindGroup: GPUBindGroup, resultBuffer: GPUBuffer) => { + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: target.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(vertexPipeline); + pass.setBindGroup(0, bindGroup); + pass.draw(kMipGradientSteps + 1); + pass.end(); + } + ), + }; device.queue.submit([encoder.finish()]); - await resultBuffer.mapAsync(GPUMapMode.READ); - const result = Array.from(new Float32Array(resultBuffer.getMappedRange())); - resultBuffer.unmap(); + await Promise.all(Object.values(stageBuffers).map(b => b.mapAsync(GPUMapMode.READ))); - const [sampleLevelWeights, gradWeights] = unzip(result, kNumWeightTypes); + const mixWeightsByStage = Object.fromEntries( + Object.entries(stageBuffers).map(([stage, resultBuffer]) => { + const result = Array.from(new Float32Array(resultBuffer.getMappedRange())); + resultBuffer.unmap(); + resultBuffer.destroy(); - validateWeights(sampleLevelWeights); - validateWeights(gradWeights); + const [sampleLevelWeights, gradWeights] = unzip(result, kNumWeightTypes); - s_deviceToMipGradientValues.set(device, { - sampleLevelWeights, - softwareMixToGPUMixGradWeights: generateSoftwareMixToGPUMixGradWeights(gradWeights, texture), - }); + validateWeights(stage, sampleLevelWeights); + validateWeights(stage, gradWeights); + + return [ + stage, + { + sampleLevelWeights, + softwareMixToGPUMixGradWeights: generateSoftwareMixToGPUMixGradWeights( + gradWeights, + texture + ), + }, + ]; + }) + ) as Record; + + s_deviceToMipGradientValues.set(device, mixWeightsByStage); texture.destroy(); storageBuffer.destroy(); - resultBuffer.destroy(); } // Given an array of ascending values and a value v, finds @@ -511,10 +719,15 @@ function generateSoftwareMixToGPUMixGradWeights(gpuWeights: number[], texture: G return softwareMixToGPUMixMap; } -function mapSoftwareMipLevelToGPUMipLevel(t: GPUTest, mipLevel: number) { +function mapSoftwareMipLevelToGPUMipLevel(t: GPUTest, stage: ShaderStage, mipLevel: number) { const baseLevel = Math.floor(mipLevel); const softwareMix = mipLevel - baseLevel; - const gpuMix = getMixWeightByTypeForMipLevel(t, 'softwareMixToGPUMixGradWeights', softwareMix); + const gpuMix = getMixWeightByTypeForMipLevel( + t, + stage, + 'softwareMixToGPUMixGradWeights', + softwareMix + ); return baseLevel + gpuMix; } @@ -549,7 +762,7 @@ const euclideanModulo = (n: number, m: number) => ((n % m) + m) % m; */ const kMipGradientSteps = 64; const s_deviceToMipGradientValuesPromise = new WeakMap>(); -const s_deviceToMipGradientValues = new WeakMap(); +const s_deviceToMipGradientValues = new WeakMap>(); async function initMipGradientValuesForDevice(t: GPUTest) { const { device } = t; let weightsP = s_deviceToMipGradientValuesPromise.get(device); @@ -562,6 +775,7 @@ async function initMipGradientValuesForDevice(t: GPUTest) { function getMixWeightByTypeForMipLevel( t: GPUTest, + stage: ShaderStage, weightType: MipWeightType | 'identity', mipLevel: number ) { @@ -569,7 +783,7 @@ function getMixWeightByTypeForMipLevel( return euclideanModulo(mipLevel, 1); } // linear interpolate between weights - const weights = s_deviceToMipGradientValues.get(t.device)![weightType]; + const weights = s_deviceToMipGradientValues.get(t.device)![stage][weightType]; assert( !!weights, 'you must use WGSLTextureSampleTest or call initializeDeviceMipWeights before calling this function' @@ -584,6 +798,7 @@ function getMixWeightByTypeForMipLevel( function getWeightForMipLevel( t: GPUTest, + stage: ShaderStage, weightType: MipWeightType | 'identity', mipLevelCount: number, mipLevel: number @@ -591,7 +806,7 @@ function getWeightForMipLevel( if (mipLevel < 0 || mipLevel >= mipLevelCount) { return 1; } - return getMixWeightByTypeForMipLevel(t, weightType, mipLevel); + return getMixWeightByTypeForMipLevel(t, stage, weightType, mipLevel); } /** @@ -1435,6 +1650,7 @@ function softwareTextureReadMipLevel( */ function softwareTextureReadLevel( t: GPUTest, + stage: ShaderStage, call: TextureCall, texture: Texture, sampler: GPUSamplerDescriptor | undefined, @@ -1456,7 +1672,7 @@ function softwareTextureReadLevel( const t0 = softwareTextureReadMipLevel(call, texture, sampler, baseMipLevel); const t1 = softwareTextureReadMipLevel(call, texture, sampler, nextMipLevel); const weightType = call.builtin === 'textureSampleLevel' ? 'sampleLevelWeights' : 'identity'; - const mix = getWeightForMipLevel(t, weightType, mipLevelCount, clampedMipLevel); + const mix = getWeightForMipLevel(t, stage, weightType, mipLevelCount, clampedMipLevel); assert(mix >= 0 && mix <= 1); const values = [ { v: t0, weight: 1 - mix }, @@ -1518,6 +1734,7 @@ function computeMipLevelFromGradientsForCall( */ function softwareTextureReadGrad( t: GPUTest, + stage: ShaderStage, call: TextureCall, texture: Texture, sampler?: GPUSamplerDescriptor @@ -1525,10 +1742,10 @@ function softwareTextureReadGrad( const bias = call.bias === undefined ? 0 : clamp(call.bias, { min: -16.0, max: 15.99 }); if (call.ddx) { const mipLevel = computeMipLevelFromGradientsForCall(call, texture.descriptor.size); - const weightMipLevel = mapSoftwareMipLevelToGPUMipLevel(t, mipLevel + bias); - return softwareTextureReadLevel(t, call, texture, sampler, weightMipLevel); + const weightMipLevel = mapSoftwareMipLevelToGPUMipLevel(t, stage, mipLevel + bias); + return softwareTextureReadLevel(t, stage, call, texture, sampler, weightMipLevel); } else { - return softwareTextureReadLevel(t, call, texture, sampler, (call.mipLevel ?? 0) + bias); + return softwareTextureReadLevel(t, stage, call, texture, sampler, (call.mipLevel ?? 0) + bias); } } @@ -1599,6 +1816,7 @@ function derivativeForCall( function softwareTextureRead( t: GPUTest, + stage: ShaderStage, call: TextureCall, texture: Texture, sampler?: GPUSamplerDescriptor @@ -1612,7 +1830,7 @@ function softwareTextureRead( }; call = newCall; } - return softwareTextureReadGrad(t, call, texture, sampler); + return softwareTextureReadGrad(t, stage, call, texture, sampler); } export type TextureTestOptions = { @@ -1841,7 +2059,8 @@ export async function checkCallResults( textureType: string, sampler: GPUSamplerDescriptor | undefined, calls: TextureCall[], - results: Awaited>> + results: Awaited>>, + stage: ShaderStage = 'fragment' // MAINTENANCE_TODO: remove default ) { const errs: string[] = []; const format = texture.texels[0].format; @@ -1856,7 +2075,7 @@ export async function checkCallResults( for (let callIdx = 0; callIdx < calls.length; callIdx++) { const call = calls[callIdx]; const gotRGBA = results.results[callIdx]; - const expectRGBA = softwareTextureRead(t, call, texture, sampler); + const expectRGBA = softwareTextureRead(t, stage, call, texture, sampler); // The spec says depth and stencil have implementation defined values for G, B, and A // so if this is `textureGather` and component > 0 then there's nothing to check. @@ -1975,6 +2194,7 @@ export async function checkCallResults( return Promise.resolve( softwareTextureRead( t, + stage, call, { texels, @@ -3863,7 +4083,10 @@ function describeTextureCall(call: TextureCall): st return `${call.builtin}(${args.join(', ')})`; } -const s_deviceToPipelines = new WeakMap>(); +const s_deviceToPipelines = new WeakMap< + GPUDevice, + Map +>(); /** * Given a list of "calls", each one of which has a texture coordinate, @@ -3937,7 +4160,8 @@ export async function doTextureCalls( viewDescriptor: GPUTextureViewDescriptor, textureType: string, sampler: GPUSamplerDescriptor | undefined, - calls: TextureCall[] + calls: TextureCall[], + stage: ShaderStage = 'fragment' // MAINTENANCE_TODO: remove default ) { const { format, @@ -4021,6 +4245,51 @@ export async function doTextureCalls( ? 'vec3f(v.pos.xy - 0.5, 0) / vec3f(textureDimensions(T))' : '(v.pos.xy - 0.5) / vec2f(textureDimensions(T))' };`; + const derivativeType = + isCubeViewDimension(viewDescriptor) || dimension === '3d' + ? 'vec3f' + : dimension === '1d' + ? 'f32' + : 'vec2f'; + + const stageWGSL = + stage === 'vertex' + ? ` +// --------------------------- vertex stage shaders -------------------------------- +@vertex fn vsVertex( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, getResult(instance_index, ${derivativeType}(0))); +} + +@fragment fn fsVertex(v: VOut) -> @location(0) vec4f { + results[v.ndx] = v.result; + return vec4f(0); +} +` + : stage === 'fragment' + ? ` +// --------------------------- fragment stage shaders -------------------------------- +@vertex fn vsFragment( + @builtin(vertex_index) vertex_index : u32, + @builtin(instance_index) instance_index : u32) -> VOut { + let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); + return VOut(vec4f(positions[vertex_index], 0, 1), instance_index, ${returnType}(0)); +} + +@fragment fn fsFragment(v: VOut) -> @location(0) vec4f { + ${derivativeBaseWGSL} + results[v.ndx] = getResult(v.ndx, derivativeBase); + return vec4f(0); +} +` + : ` +// --------------------------- compute stage shaders -------------------------------- +@compute @workgroup_size(1) fn csCompute(@builtin(global_invocation_id) id: vec3u) { + results[id.x] = getResult(id.x, ${derivativeType}(0)); +} +`; const code = ` ${structs} @@ -4032,33 +4301,25 @@ ${dataFields} struct VOut { @builtin(position) pos: vec4f, @location(0) @interpolate(flat, either) ndx: u32, + @location(1) @interpolate(flat, either) result: ${returnType}, }; -@vertex -fn vs_main( - @builtin(vertex_index) vertex_index : u32, - @builtin(instance_index) instance_index : u32) -> VOut { - let positions = array(vec2f(-1, 3), vec2f(3, -1), vec2f(-1, -1)); - return VOut(vec4f(positions[vertex_index], 0, 1), instance_index); -} - @group(0) @binding(0) var T : ${textureType}; ${sampler ? `@group(0) @binding(1) var S : ${samplerType}` : ''}; @group(0) @binding(2) var data : Data; -@group(0) @binding(3) var results: array<${returnType}>; +@group(1) @binding(0) var results: array<${returnType}>; -@fragment -fn fs_main(v: VOut) -> @location(0) vec4f { - ${derivativeBaseWGSL} - let idx = v.ndx; +fn getResult(idx: u32, derivativeBase: ${derivativeType}) -> ${returnType} { var result : ${resultType}; ${body} - results[idx] = ${returnType}(result); - return vec4f(0); + return ${returnType}(result); } + +${stageWGSL} `; - const pipelines = s_deviceToPipelines.get(t.device) ?? new Map(); + const pipelines = + s_deviceToPipelines.get(t.device) ?? new Map(); s_deviceToPipelines.set(t.device, pipelines); // unfilterable-float textures can only be used with manually created bindGroupLayouts @@ -4088,21 +4349,21 @@ ${body} sampleType = 'unfilterable-float'; } + const visibility = + stage === 'compute' + ? GPUShaderStage.COMPUTE + : stage === 'fragment' + ? GPUShaderStage.FRAGMENT + : GPUShaderStage.VERTEX; + const entries: GPUBindGroupLayoutEntry[] = [ { binding: 2, - visibility: GPUShaderStage.FRAGMENT, + visibility, buffer: { type: 'read-only-storage', }, }, - { - binding: 3, - visibility: GPUShaderStage.FRAGMENT, - buffer: { - type: 'storage', - }, - }, ]; const viewDimension = effectiveViewDimensionForDimension( @@ -4114,7 +4375,7 @@ ${body} if (textureType.includes('storage')) { entries.push({ binding: 0, - visibility: GPUShaderStage.FRAGMENT, + visibility, storageTexture: { access: 'read-only', viewDimension, @@ -4124,13 +4385,13 @@ ${body} } else if (gpuTexture instanceof GPUExternalTexture) { entries.push({ binding: 0, - visibility: GPUShaderStage.FRAGMENT, + visibility, externalTexture: {}, }); } else { entries.push({ binding: 0, - visibility: GPUShaderStage.FRAGMENT, + visibility, texture: { sampleType, viewDimension, @@ -4142,31 +4403,64 @@ ${body} if (sampler) { entries.push({ binding: 1, - visibility: GPUShaderStage.FRAGMENT, + visibility, sampler: { type: isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering', }, }); } - const id = `${resultType}:${JSON.stringify(entries)}:${code}`; + const id = `${resultType}:${stage}:${JSON.stringify(entries)}:${code}`; let pipeline = pipelines.get(id); if (!pipeline) { - const shaderModule = t.device.createShaderModule({ code }); - const bindGroupLayout = t.device.createBindGroupLayout({ entries }); - const layout = t.device.createPipelineLayout({ - bindGroupLayouts: [bindGroupLayout], + const module = t.device.createShaderModule({ code }); + const bindGroupLayout0 = t.device.createBindGroupLayout({ entries }); + const bindGroupLayout1 = t.device.createBindGroupLayout({ + entries: [ + { + binding: 0, + visibility: GPUShaderStage.FRAGMENT | GPUShaderStage.COMPUTE, + buffer: { + type: 'storage', + }, + }, + ], }); - pipeline = t.device.createRenderPipeline({ - layout, - vertex: { module: shaderModule }, - fragment: { - module: shaderModule, - targets: [{ format: renderTarget.format }], - }, + const layout = t.device.createPipelineLayout({ + bindGroupLayouts: [bindGroupLayout0, bindGroupLayout1], }); + switch (stage) { + case 'compute': + pipeline = t.device.createComputePipeline({ + layout, + compute: { module }, + }); + break; + case 'fragment': + pipeline = t.device.createRenderPipeline({ + layout, + vertex: { module, entryPoint: 'vsFragment' }, + fragment: { + module, + entryPoint: 'fsFragment', + targets: [{ format: renderTarget.format }], + }, + }); + break; + case 'vertex': + pipeline = t.device.createRenderPipeline({ + layout, + vertex: { module, entryPoint: 'vsVertex' }, + fragment: { + module, + entryPoint: 'fsVertex', + targets: [{ format: renderTarget.format }], + }, + }); + break; + } pipelines.set(id, pipeline); } @@ -4178,7 +4472,7 @@ ${body} usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); - const bindGroup = t.device.createBindGroup({ + const bindGroup0 = t.device.createBindGroup({ layout: pipeline!.getBindGroupLayout(0), entries: [ { @@ -4190,10 +4484,14 @@ ${body} }, ...(sampler ? [{ binding: 1, resource: gpuSampler! }] : []), { binding: 2, resource: { buffer: dataBuffer } }, - { binding: 3, resource: { buffer: storageBuffer } }, ], }); + const bindGroup1 = t.device.createBindGroup({ + layout: pipeline!.getBindGroupLayout(1), + entries: [{ binding: 0, resource: { buffer: storageBuffer } }], + }); + const resultBuffer = t.createBufferTracked({ size: storageBuffer.size, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, @@ -4201,20 +4499,30 @@ ${body} const encoder = t.device.createCommandEncoder(); - const renderPass = encoder.beginRenderPass({ - colorAttachments: [ - { - view: renderTarget.createView(), - loadOp: 'clear', - storeOp: 'store', - }, - ], - }); + if (stage === 'compute') { + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline! as GPUComputePipeline); + pass.setBindGroup(0, bindGroup0); + pass.setBindGroup(1, bindGroup1); + pass.dispatchWorkgroups(calls.length); + pass.end(); + } else { + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: renderTarget.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); - renderPass.setPipeline(pipeline!); - renderPass.setBindGroup(0, bindGroup); - renderPass.draw(3, calls.length); - renderPass.end(); + pass.setPipeline(pipeline! as GPURenderPipeline); + pass.setBindGroup(0, bindGroup0); + pass.setBindGroup(1, bindGroup1); + pass.draw(3, calls.length); + pass.end(); + } encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, storageBuffer.size); t.device.queue.submit([encoder.finish()]);