From d849b67718d5166f3ab9479c91570a612ad2edd7 Mon Sep 17 00:00:00 2001 From: David Neto Date: Tue, 13 Aug 2024 15:05:50 -0400 Subject: [PATCH 01/99] validation: clamp low_high test must reference the test function (#3899) The clamp validation test creates a 'foo' function that contains code to be checked. The entry point must call 'foo' in order for those override expressions to be checked. Bug: crbug.com/351378281 --- src/webgpu/listing_meta.json | 2 +- .../shader/validation/expression/call/builtin/clamp.spec.ts | 1 + src/webgpu/shader/validation/shader_validation_test.ts | 5 +++++ 3 files changed, 7 insertions(+), 1 deletion(-) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 85fe0bdc6a8b..2374c4cff7e1 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -862,7 +862,6 @@ "webgpu:compat,api,validation,encoding,programmable,pipeline_bind_group_compat:twoDifferentTextureViews,render_pass,used:*": { "subcaseMS": 0.000 }, "webgpu:compat,api,validation,render_pipeline,depth_stencil_state:depthBiasClamp:*": { "subcaseMS": 1.604 }, "webgpu:compat,api,validation,render_pipeline,fragment_state:colorState:*": { "subcaseMS": 32.604 }, - "webgpu:compat,api,validation,render_pipeline,vertex_state:maxVertexAttributesVertexIndexInstanceIndex:*": { "subcaseMS": 3.700 }, "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:interpolate:*": { "subcaseMS": 3.488 }, "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:sample_index:*": { "subcaseMS": 0.487 }, "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:sample_mask:*": { "subcaseMS": 0.408 }, @@ -870,6 +869,7 @@ "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:textureLoad_with_depth_textures,renderPipeline:*": { "subcaseMS": 1.259 }, "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:unsupportedStorageTextureFormats,computePipeline:*": { "subcaseMS": 1.206 }, "webgpu:compat,api,validation,render_pipeline,unsupported_wgsl:unsupportedStorageTextureFormats,renderPipeline:*": { "subcaseMS": 1.206 }, + "webgpu:compat,api,validation,render_pipeline,vertex_state:maxVertexAttributesVertexIndexInstanceIndex:*": { "subcaseMS": 3.700 }, "webgpu:compat,api,validation,texture,createTexture:depthOrArrayLayers_incompatible_with_textureBindingViewDimension:*": { "subcaseMS": 12.712 }, "webgpu:compat,api,validation,texture,createTexture:format_reinterpretation:*": { "subcaseMS": 7.012 }, "webgpu:compat,api,validation,texture,createTexture:invalidTextureBindingViewDimension:*": { "subcaseMS": 6.022 }, diff --git a/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts index 1ac752a3bfa9..e94162b1ce5a 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts @@ -187,6 +187,7 @@ fn foo() { code: wgsl, constants, reference: ['o_low', 'o_high'], + statements: ['foo();'], }); } }); diff --git a/src/webgpu/shader/validation/shader_validation_test.ts b/src/webgpu/shader/validation/shader_validation_test.ts index 6a4cae331766..5db47bd586ba 100644 --- a/src/webgpu/shader/validation/shader_validation_test.ts +++ b/src/webgpu/shader/validation/shader_validation_test.ts @@ -119,9 +119,14 @@ export class ShaderValidationTest extends GPUTest { constants?: Record; // List of additional module-scope variable the entrypoint needs to reference reference?: string[]; + // List of additional statements to insert in the entry point. + statements?: string[]; }) { const phonies: Array = []; + if (args.statements !== undefined) { + phonies.push(...args.statements); + } if (args.constants !== undefined) { phonies.push(...keysOf(args.constants).map(c => `_ = ${c};`)); } From 24725f3c305f0b3f1e3a8add5a262871364a9030 Mon Sep 17 00:00:00 2001 From: David Neto Date: Tue, 13 Aug 2024 17:31:07 -0400 Subject: [PATCH 02/99] fix smoothstep validation (#3900) 'values' test: * Infer result type instead of explicitly specifying it. If specified, it's used as the explicit type of the result variable. Infer it instead. This fixes cases where the type is abstract. * low == high is an error https://github.com/gpuweb/gpuweb/pull/4616 'partial_eval_errs' test: * Test function foo() must be called from the entry point so that overrides are validated. 'early_eval_errs' test: * Fix argument order: low and high args are first and second. Bug: crbug.com/351378281 --- .../expression/call/builtin/smoothstep.spec.ts | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts index 5a5a28fc7362..d9dd63b85bf3 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts @@ -51,16 +51,15 @@ Validates that constant evaluation and override evaluation of ${builtin}() rejec .fn(t => { const type = kValuesTypes[t.params.type]; - // We expect to fail if low >= high as it results in a DBZ - const expectedResult = t.params.value1 >= t.params.value2; + // We expect to fail if low >= high. + const expectedResult = t.params.value1 < t.params.value2; validateConstOrOverrideBuiltinEval( t, builtin, expectedResult, [type.create(t.params.value1), type.create(t.params.value2), type.create(0)], - t.params.stage, - /* returnType */ concreteTypeOf(type, [Type.f32]) + t.params.stage ); }); @@ -141,6 +140,7 @@ fn foo() { code: wgsl, constants, reference: ['o_low', 'o_high'], + statements: ['foo();'], }); } }); @@ -159,10 +159,11 @@ Validates that scalar and vector arguments are rejected by ${builtin}() if not f }) .fn(t => { const type = kArgumentTypes[t.params.type]; + const expectedResult = isConvertibleToFloatType(elementTypeOf(type)); validateConstOrOverrideBuiltinEval( t, builtin, - /* expectedResult */ isConvertibleToFloatType(elementTypeOf(type)), + expectedResult, [type.create(0), type.create(1), type.create(2)], 'constant', /* returnType */ concreteTypeOf(type, [Type.f32]) @@ -344,7 +345,7 @@ g.test('early_eval_errors') t, builtin, /* expectedResult */ t.params.low < t.params.high, - [f32(0), f32(t.params.low), f32(t.params.high)], + [f32(t.params.low), f32(t.params.high), f32(0)], t.params.stage ); }); From 701e646366987461e8019725a26ee527f7f8c735 Mon Sep 17 00:00:00 2001 From: Greggman Date: Wed, 14 Aug 2024 08:46:44 -0700 Subject: [PATCH 03/99] Fix textureSampleLevel tests. (#3903) The tests query the GPUs mip to mip interpolation. The test that the query was successful was too strict so this PR relaxes that check. --- .../execution/expression/call/builtin/texture_utils.ts | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) 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 e997833a137f..38e94675f3e0 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -203,9 +203,12 @@ async function initMipGradientValuesForDevice(t: GPUTest) { resultBuffer.destroy(); // Validate the weights - assert(weights[0] === 0); - assert(weights[kMipGradientSteps] === 1); - assert(weights[kMipGradientSteps / 2] === 0.5); + assert(weights[0] === 0, 'weight 0 is 0'); + assert(weights[kMipGradientSteps] === 1, 'top weight is 1'); + assert( + Math.abs(weights[kMipGradientSteps / 2] - 0.5) < 0.0001, + 'middle weight is approximately 0.5' + ); // Note: for 16 steps, these are the AMD weights // From a17d71da10b513be9f42342337f1ac07d9f17431 Mon Sep 17 00:00:00 2001 From: David Neto Date: Wed, 14 Aug 2024 12:45:57 -0400 Subject: [PATCH 04/99] fix smoothstep execution: only compute valid cases when const (#3901) For const cases, low < high is required See https://github.com/gpuweb/gpuweb/pull/4616 Bug: crbug.com/351378281 --- .../call/builtin/smoothstep.spec.ts | 36 ++++++++++++++++--- .../call/builtin/smoothstep.spec.ts | 9 +++-- 2 files changed, 39 insertions(+), 6 deletions(-) diff --git a/src/webgpu/shader/execution/expression/call/builtin/smoothstep.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/smoothstep.spec.ts index 42d8d09ff569..f65bb951bf25 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/smoothstep.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/smoothstep.spec.ts @@ -7,11 +7,16 @@ T is S or vecN Returns the smooth Hermite interpolation between 0 and 1. Component-wise when T is a vector. For scalar T, the result is t * t * (3.0 - 2.0 * t), where t = clamp((x - low) / (high - low), 0.0, 1.0). + +If low >= high: +* It is a shader-creation error if low and high are const-expressions. +* It is a pipeline-creation error if low and high are override-expressions. `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; import { GPUTest } from '../../../../../gpu_test.js'; -import { Type } from '../../../../../util/conversion.js'; +import { ScalarValue, Type, Value } from '../../../../../util/conversion.js'; +import { Case } from '../../case.js'; import { allInputSources, onlyConstInputSource, run } from '../../expression.js'; import { abstractFloatBuiltin, builtin } from './builtin.js'; @@ -19,6 +24,13 @@ import { d } from './smoothstep.cache.js'; export const g = makeTestGroup(GPUTest); +// Returns true if `c` is valid for a const evaluation of smoothstep. +function validForConst(c: Case): boolean { + const low = (c.input as Value[])[0] as ScalarValue; + const high = (c.input as Value[])[1] as ScalarValue; + return low.value < high.value; +} + g.test('abstract_float') .specURL('https://www.w3.org/TR/WGSL/#float-builtin-functions') .desc(`abstract float tests`) @@ -28,7 +40,7 @@ g.test('abstract_float') .combine('vectorize', [undefined, 2, 3, 4] as const) ) .fn(async t => { - const cases = await d.get('abstract_const'); + const cases = (await d.get('abstract_const')).filter(c => validForConst(c)); await run( t, abstractFloatBuiltin('smoothstep'), @@ -47,7 +59,15 @@ g.test('f32') ) .fn(async t => { const cases = await d.get(t.params.inputSource === 'const' ? 'f32_const' : 'f32_non_const'); - await run(t, builtin('smoothstep'), [Type.f32, Type.f32, Type.f32], Type.f32, t.params, cases); + const validCases = cases.filter(c => t.params.inputSource !== 'const' || validForConst(c)); + await run( + t, + builtin('smoothstep'), + [Type.f32, Type.f32, Type.f32], + Type.f32, + t.params, + validCases + ); }); g.test('f16') @@ -61,5 +81,13 @@ g.test('f16') }) .fn(async t => { const cases = await d.get(t.params.inputSource === 'const' ? 'f16_const' : 'f16_non_const'); - await run(t, builtin('smoothstep'), [Type.f16, Type.f16, Type.f16], Type.f16, t.params, cases); + const validCases = cases.filter(c => t.params.inputSource !== 'const' || validForConst(c)); + await run( + t, + builtin('smoothstep'), + [Type.f16, Type.f16, Type.f16], + Type.f16, + t.params, + validCases + ); }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts index d9dd63b85bf3..2879055ab216 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/smoothstep.spec.ts @@ -80,6 +80,8 @@ g.test('partial_eval_errors') .beginSubcases() .expand('low', u => [0, 10]) .expand('high', u => [0, 10]) + // in_shader: Is the function call statically accessed by the entry point? + .combine('in_shader', [false, true] as const) ) .beforeAllSubcases(t => { if (scalarTypeOf(kValuesTypes[t.params.type]) === Type.f16) { @@ -129,7 +131,10 @@ fn foo() { const shader_error = error && t.params.lowStage === 'constant' && t.params.highStage === 'constant'; const pipeline_error = - error && t.params.lowStage !== 'runtime' && t.params.highStage !== 'runtime'; + t.params.in_shader && + error && + t.params.lowStage !== 'runtime' && + t.params.highStage !== 'runtime'; t.expectCompileResult(!shader_error, wgsl); if (!shader_error) { const constants: Record = {}; @@ -140,7 +145,7 @@ fn foo() { code: wgsl, constants, reference: ['o_low', 'o_high'], - statements: ['foo();'], + statements: t.params.in_shader ? ['foo();'] : [], }); } }); From 5c571490b4c077f0319cb6a9b06d77ed39e205dd Mon Sep 17 00:00:00 2001 From: David Neto Date: Wed, 14 Aug 2024 15:38:20 -0400 Subject: [PATCH 05/99] clamp validation: conditionally include test function (#3904) Add variants where the function containing the function call being tested is both in and not-in the tested shader, i.e. statically accessed or not. This matters for override validation. --- .../validation/expression/call/builtin/clamp.spec.ts | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts index e94162b1ce5a..ff0114097f90 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/clamp.spec.ts @@ -127,6 +127,8 @@ Validates that low <= high. const scalar = scalarTypeOf(ty); return scalar !== Type.abstractInt && scalar !== Type.abstractFloat; }) + // in_shader: Is the function call statically accessed by the entry point? + .combine('in_shader', [false, true] as const) ) .beforeAllSubcases(t => { const ty = kValuesTypes[t.params.type]; @@ -176,7 +178,10 @@ fn foo() { const shader_error = error && t.params.lowStage === 'constant' && t.params.highStage === 'constant'; const pipeline_error = - error && t.params.lowStage !== 'runtime' && t.params.highStage !== 'runtime'; + t.params.in_shader && + error && + t.params.lowStage !== 'runtime' && + t.params.highStage !== 'runtime'; t.expectCompileResult(!shader_error, wgsl); if (!shader_error) { const constants: Record = {}; @@ -187,7 +192,7 @@ fn foo() { code: wgsl, constants, reference: ['o_low', 'o_high'], - statements: ['foo();'], + statements: t.params.in_shader ? ['foo();'] : [], }); } }); From 67dc033599d4292dcacca587bac78855ac96d8d0 Mon Sep 17 00:00:00 2001 From: Greggman Date: Wed, 14 Aug 2024 23:06:18 -0700 Subject: [PATCH 06/99] Compat: Fix textureLoad with depth texture test (#3902) --- .../api/validation/render_pipeline/unsupported_wgsl.spec.ts | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/webgpu/compat/api/validation/render_pipeline/unsupported_wgsl.spec.ts b/src/webgpu/compat/api/validation/render_pipeline/unsupported_wgsl.spec.ts index 805203870903..27b25f34af49 100644 --- a/src/webgpu/compat/api/validation/render_pipeline/unsupported_wgsl.spec.ts +++ b/src/webgpu/compat/api/validation/render_pipeline/unsupported_wgsl.spec.ts @@ -267,7 +267,7 @@ g.test('textureLoad_with_depth_textures,computePipeline') `, }); - const isValid = !t.isCompatibility; + const isValid = !t.isCompatibility || entryPoint === 'csWithoutDepthUsage'; t.doCreateComputePipelineTest(async, isValid, { layout: 'auto', compute: { module, entryPoint }, @@ -301,7 +301,7 @@ g.test('textureLoad_with_depth_textures,renderPipeline') `, }); - const isValid = !t.isCompatibility; + const isValid = !t.isCompatibility || entryPoint === 'vsWithoutDepthUsage'; t.doCreateRenderPipelineTest(async, isValid, { layout: 'auto', vertex: { module, entryPoint }, From c279d024b0934cf2ed5e32b9ff215d580bedfaba Mon Sep 17 00:00:00 2001 From: David Neto Date: Thu, 15 Aug 2024 10:49:37 -0400 Subject: [PATCH 07/99] extractBits validation: conditionally include tested code in the shader (#3906) This matters when exercising the'override' cases. --- .../expression/call/builtin/extractBits.spec.ts | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/src/webgpu/shader/validation/expression/call/builtin/extractBits.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/extractBits.spec.ts index 80fe7ccaca5e..32abc477ee8f 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/extractBits.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/extractBits.spec.ts @@ -98,6 +98,8 @@ Validates that count and offset must be smaller than the size of the primitive. { offset: 0, count: 33 }, { offset: 1, count: 33 }, ] as const) + // in_shader: Is the function call statically accessed by the entry point? + .combine('in_shader', [false, true] as const) ) .fn(t => { let offsetArg = ''; @@ -138,7 +140,10 @@ fn foo() { const shader_error = error && t.params.offsetStage === 'constant' && t.params.countStage === 'constant'; const pipeline_error = - error && t.params.offsetStage !== 'runtime' && t.params.countStage !== 'runtime'; + t.params.in_shader && + error && + t.params.offsetStage !== 'runtime' && + t.params.countStage !== 'runtime'; t.expectCompileResult(!shader_error, wgsl); if (!shader_error) { const constants: Record = {}; @@ -149,6 +154,7 @@ fn foo() { code: wgsl, constants, reference: ['o_offset', 'o_count'], + statements: t.params.in_shader ? ['foo();'] : [], }); } }); From d4d5926657beddb7c3ea70016942ce46a45467cf Mon Sep 17 00:00:00 2001 From: David Neto Date: Thu, 15 Aug 2024 12:47:56 -0400 Subject: [PATCH 08/99] insertBits validation: conditionally include tested code in the shader (#3905) This matters when exercising the 'override' cases. --- .../validation/expression/call/builtin/insertBits.spec.ts | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/src/webgpu/shader/validation/expression/call/builtin/insertBits.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/insertBits.spec.ts index 57644ad36fb4..b302bfd14677 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/insertBits.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/insertBits.spec.ts @@ -119,6 +119,8 @@ Validates that count and offset must be smaller than the size of the primitive. { offset: 0, count: 33 }, { offset: 1, count: 33 }, ] as const) + // in_shader: Is the function call statically accessed by the entry point? + .combine('in_shader', [false, true] as const) ) .fn(t => { let offsetArg = ''; @@ -160,7 +162,10 @@ fn foo() { const shader_error = error && t.params.offsetStage === 'constant' && t.params.countStage === 'constant'; const pipeline_error = - error && t.params.offsetStage !== 'runtime' && t.params.countStage !== 'runtime'; + t.params.in_shader && + error && + t.params.offsetStage !== 'runtime' && + t.params.countStage !== 'runtime'; t.expectCompileResult(!shader_error, wgsl); if (!shader_error) { const constants: Record = {}; @@ -171,6 +176,7 @@ fn foo() { code: wgsl, constants, reference: ['o_offset', 'o_count'], + statements: t.params.in_shader ? ['foo();'] : [], }); } }); From 4ed3eaf11314f23289195346f5849c21ce56c18e Mon Sep 17 00:00:00 2001 From: Greggman Date: Thu, 15 Aug 2024 10:44:43 -0700 Subject: [PATCH 09/99] Fix for texture_utils.ts (#3907) The issue is on Chrome and Firefox on Intel Mac, when sampling between 2 mip levels using `textureSampleLevel` in a compute shader, the weights used for mixing are very unexpected. The same issue doesn't happen in Safari TP so there is probably a fix. For now though, the same issue doesn't happen when using a fragment shader. So, switched to using a fragment shader to look up these weights. This is more appropriate for the current tests because the tests are running in fragment shaders. Will add an issue to test all stages. --- .../expression/call/builtin/texture_utils.ts | 40 +++++++++++++++---- 1 file changed, 32 insertions(+), 8 deletions(-) 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 38e94675f3e0..7e8b1e168482 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -134,16 +134,32 @@ async function initMipGradientValuesForDevice(t: GPUTest) { @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; + @vertex fn vs(@builtin(vertex_index) vNdx: u32) -> @builtin(position) vec4f { + let pos = array( + vec2f(-1, 3), + vec2f( 3, -1), + vec2f(-1, -1), + ); + return vec4f(pos[vNdx], 0, 1); + } + @fragment fn fs(@builtin(position) pos: vec4f) -> @location(0) vec4f { + let mipLevel = floor(pos.x) / ${kMipGradientSteps}; + result[u32(pos.x)] = textureSampleLevel(tex, smp, vec2f(0.5), mipLevel).r; + return vec4f(0); } `, }); - const pipeline = device.createComputePipeline({ + const pipeline = device.createRenderPipeline({ layout: 'auto', - compute: { module }, + vertex: { module }, + fragment: { module, targets: [{ format: 'rgba8unorm' }] }, + }); + + const target = t.createTextureTracked({ + size: [kMipGradientSteps + 1, 1, 1], + format: 'rgba8unorm', + usage: GPUTextureUsage.RENDER_ATTACHMENT, }); const texture = t.createTextureTracked({ @@ -186,10 +202,18 @@ async function initMipGradientValuesForDevice(t: GPUTest) { }); const encoder = device.createCommandEncoder(); - const pass = encoder.beginComputePass(); + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: target.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(kMipGradientSteps + 1); + pass.draw(3); pass.end(); encoder.copyBufferToBuffer(storageBuffer, 0, resultBuffer, 0, resultBuffer.size); device.queue.submit([encoder.finish()]); @@ -215,7 +239,7 @@ async function initMipGradientValuesForDevice(t: GPUTest) { // standard // step mipLevel gpu AMD // ---- -------- -------- ---------- - // 0: 0 0 1 + // 0: 0 0 0 // 1: 0.0625 0.0625 0 // 2: 0.125 0.125 0.03125 // 3: 0.1875 0.1875 0.109375 From 4ccb5a4736f18db85582438d1c27112609f13c9e Mon Sep 17 00:00:00 2001 From: Greggman Date: Thu, 15 Aug 2024 13:14:04 -0700 Subject: [PATCH 10/99] Stop Sooner (#3889) I noticed when I click stop, sometimes it seems ignored. I think this is why. --- src/common/runtime/standalone.ts | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/common/runtime/standalone.ts b/src/common/runtime/standalone.ts index 932c5668b587..0305031cc790 100644 --- a/src/common/runtime/standalone.ts +++ b/src/common/runtime/standalone.ts @@ -369,6 +369,9 @@ function makeSubtreeChildrenHTML( const runMySubtree = async () => { const results: SubtreeResult[] = []; for (const { runSubtree } of childFns) { + if (stopRequested) { + break; + } results.push(await runSubtree()); } return mergeSubtreeResults(...results); From a96693cd274c5b340ffe37df6ff265596059dd1e Mon Sep 17 00:00:00 2001 From: David Neto Date: Fri, 16 Aug 2024 18:44:56 -0400 Subject: [PATCH 11/99] ldexp validation: conditionally include tested code in the shader (#3908) This affects 'override' cases in the 'partial_values' subtest. --- .../shader/validation/expression/call/builtin/ldexp.spec.ts | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/webgpu/shader/validation/expression/call/builtin/ldexp.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/ldexp.spec.ts index 826354d1ff08..55a702d71f0a 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/ldexp.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/ldexp.spec.ts @@ -143,6 +143,8 @@ g.test('partial_values') cases.push({ value: bias + 2 }); return cases; }) + // in_shader: Is the functino call statically accessed by the entry point? + .combine('in_shader', [false, true] as const) ) .beforeAllSubcases(t => { const ty = kValidArgumentTypesA[t.params.typeA]; @@ -179,7 +181,7 @@ fn foo() { const bias = biasForType(scalarTypeOf(tyA)); const error = t.params.value > bias + 1; const shader_error = error && t.params.stage === 'constant'; - const pipeline_error = error && t.params.stage === 'override'; + const pipeline_error = t.params.in_shader && error && t.params.stage === 'override'; t.expectCompileResult(!shader_error, wgsl); if (!shader_error) { const constants: Record = {}; @@ -189,6 +191,7 @@ fn foo() { code: wgsl, constants, reference: ['o_b'], + statements: t.params.in_shader ? ['foo();'] : [], }); } }); From 32559adb87c220a0d99e1b7cfafec7f835e102b3 Mon Sep 17 00:00:00 2001 From: Greggman Date: Thu, 22 Aug 2024 04:12:11 +0900 Subject: [PATCH 12/99] Reduce the number of shader modules used in texture tests (#3910) This change reduces the number of shader modules created by 2 orders of magnitude. The issue is offsets must be constants so randomly generating them makes a new shader. This makes them less random so they are the same per test. So for example: webgpu:shader,execution,expression,call,builtin,textureSampleLevel:sampled_2d_coords:* goes from 1197 shader modules to 4 --- .../execution/expression/call/builtin/texture_utils.ts | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) 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 7e8b1e168482..3029235976c7 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -2266,9 +2266,10 @@ function generateTextureBuiltinInputsImpl( (hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1); return type === 'f32' ? number : Math.floor(number); }; - const makeIntHashValue = (min: number, max: number, ...hashInputs: number[]) => { + // Generates the same values per coord instead of using all the extra `_hashInputs`. + const makeIntHashValueRepeatable = (min: number, max: number, ...hashInputs: number[]) => { const range = max - min; - return min + Math.floor((hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range); + return min + Math.floor((hashU32(...hashInputs) / 0x1_0000_0000) * range); }; // Samplers across devices use different methods to interpolate. @@ -2305,7 +2306,7 @@ function generateTextureBuiltinInputsImpl( sampleIndex: args.sampleIndex ? makeRangeValue(args.sampleIndex, i, 1) : undefined, arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined, offset: args.offset - ? (coords.map((_, j) => makeIntHashValue(-8, 8, i, 3 + j)) as T) + ? (coords.map((_, j) => makeIntHashValueRepeatable(-8, 8, i, 3 + j)) as T) : undefined, }; }); From f8472c9ab5cb878832d3e20ec074ca634215fad6 Mon Sep 17 00:00:00 2001 From: Greggman Date: Thu, 22 Aug 2024 04:22:17 +0900 Subject: [PATCH 13/99] Print more info in certain asserts (#3911) This is useful for CQs where bots have different hardware then devs have locally. --- .../expression/call/builtin/texture_utils.ts | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) 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 3029235976c7..f996c6ce1c7e 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -226,12 +226,19 @@ async function initMipGradientValuesForDevice(t: GPUTest) { storageBuffer.destroy(); resultBuffer.destroy(); + const showWeights = () => weights.map((v, i) => `${i.toString().padStart(2)}: ${v}`).join('\n'); + // Validate the weights - assert(weights[0] === 0, 'weight 0 is 0'); - assert(weights[kMipGradientSteps] === 1, 'top weight is 1'); + 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()}` + ); assert( Math.abs(weights[kMipGradientSteps / 2] - 0.5) < 0.0001, - 'middle weight is approximately 0.5' + `middle weight expected approximately 0.5 but was ${ + weights[kMipGradientSteps / 2] + }\n${showWeights()}` ); // Note: for 16 steps, these are the AMD weights From 66327462a9b5f01270f39ccf158372d31563dcd7 Mon Sep 17 00:00:00 2001 From: Greggman Date: Thu, 22 Aug 2024 05:38:48 +0900 Subject: [PATCH 14/99] Choose smaller texture sizes (#3912) The old code did an lcm of width vs height where for for cubemaps because cubemaps must be square and textures must be a multiple of block sizes. With a format with a blockSize of 5x8 and a minSize of 32 that would end up doing lcm of 35x32 which is 1120. Then for a cube array it would end up allocating 1120x1120x24 and if the format is rgba32float that's 418meg. The new code just gets the lcm of the blockWidth vs blockHeight and then aligning to that which will be much much smaller. --- .../shader/execution/expression/call/builtin/texture_utils.ts | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) 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 f996c6ce1c7e..3f1d778b52b1 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -2197,7 +2197,9 @@ export function chooseTextureSize({ const width = align(Math.max(minSize, blockWidth * minBlocks), blockWidth); const height = align(Math.max(minSize, blockHeight * minBlocks), blockHeight); if (viewDimension === 'cube' || viewDimension === 'cube-array') { - const size = lcm(width, height); + const blockLCM = lcm(blockWidth, blockHeight); + const largest = Math.max(width, height); + const size = align(largest, blockLCM); return [size, size, viewDimension === 'cube-array' ? 24 : 6]; } const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); From 348b113f26b0c5c007d37e3379cb32600a9defb4 Mon Sep 17 00:00:00 2001 From: petermcneeleychromium <96925679+petermcneeleychromium@users.noreply.github.com> Date: Thu, 22 Aug 2024 12:37:07 -0400 Subject: [PATCH 15/99] Fix for 16f min (#3913) Co-authored-by: Peter McNeeley --- src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts b/src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts index 85bed5228482..d162ba3286b9 100644 --- a/src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts +++ b/src/webgpu/shader/validation/expression/matrix/add_sub.spec.ts @@ -275,7 +275,7 @@ g.test('underflow_f16') let rhs = `mat${t.params.c}x${t.params.r}h(`; for (let i = 0; i < t.params.c; i++) { for (let k = 0; k < t.params.r; k++) { - lhs += `${kValue.f32.negative.min / 2},`; + lhs += `${kValue.f16.negative.min / 2},`; rhs += `${t.params.rhs},`; } } From b2e8800940095bc0e7bf065797257613fedb4aef Mon Sep 17 00:00:00 2001 From: Greggman Date: Fri, 23 Aug 2024 02:40:16 +0900 Subject: [PATCH 16/99] WGSL textureGather tests (#3898) --- .../call/builtin/textureGather.spec.ts | 630 ++++++++++++-- .../call/builtin/textureLoad.spec.ts | 13 +- .../call/builtin/textureSampleLevel.spec.ts | 52 +- .../expression/call/builtin/texture_utils.ts | 817 +++++++++++++----- src/webgpu/util/texture.ts | 256 +++--- 5 files changed, 1320 insertions(+), 448 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 40b331efaba9..f0f955b87dc8 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts @@ -1,6 +1,8 @@ export const description = ` Execution tests for the 'textureGather' builtin function +- TODO: Test un-encodable formats. + A texture gather operation reads from a 2D, 2D array, cube, or cube array texture, computing a four-component vector as follows: * Find the four texels that would be used in a sampling operation with linear filtering, from mip level 0: - Use the specified coordinate, array index (when present), and offset (when present). @@ -23,11 +25,35 @@ A texture gather operation reads from a 2D, 2D array, cube, or cube array textur `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; +import { + isDepthTextureFormat, + isEncodableTextureFormat, + kCompressedTextureFormats, + kDepthStencilFormats, + kEncodableTextureFormats, +} from '../../../../../format_info.js'; + +import { + appendComponentTypeForFormatToTextureType, + checkCallResults, + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + doTextureCalls, + generateSamplePointsCube, + generateTextureBuiltinInputs2D, + isFillable, + kCubeSamplePointMethods, + kSamplePointMethods, + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice, + 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/#texturegather') @@ -55,22 +81,77 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('T', ['f32-only', 'i32', 'u32'] as const) - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) + .combine('format', kTestableColorFormats) + .filter(t => isFillable(t.format)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() .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('samplePoints', kSamplePointMethods) + .combine('addressModeU', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('addressModeV', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); + }) + .fn(async t => { + const { format, C, 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, + offset, + component: true, + hashInputs: [format, C, samplePoints, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, component, offset }) => { + return { + builtin: 'textureGather', + coordType: 'f', + coords, + component, + componentType: C === 'i32' ? 'i' : 'u', + 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_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegather') .desc( ` -C: i32, u32 T: i32, u32, f32 fn textureGather(component: C, t: texture_cube, s: sampler, coords: vec3) -> vec4 @@ -85,15 +166,75 @@ Parameters: * coords: The texture coordinates ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('T', ['f32-only', 'i32', 'u32'] as const) - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) + .combine('format', kTestableColorFormats) + .filter(t => isFillable(t.format)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(3)) + .combine('samplePoints', kCubeSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); + }) + .fn(async t => { + const { format, C, samplePoints, addressMode, minFilter } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube'; + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + const depthOrArrayLayers = 6; + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size: { width, height, depthOrArrayLayers }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + mipLevelCount: 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[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + component: true, + textureBuiltin: 'textureGather', + hashInputs: [format, C, samplePoints, addressMode, minFilter], + }).map(({ coords, component }) => { + return { + builtin: 'textureGather', + component, + componentType: C === 'i32' ? 'i' : 'u', + coordType: 'f', + coords, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = appendComponentTypeForFormatToTextureType('texture_cube', format); + 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/#texturegather') @@ -122,17 +263,78 @@ Parameters: Values outside of this range will result in a shader-creation error. ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('T', ['f32-only', 'i32', 'u32'] as const) - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) + .combine('format', kTestableColorFormats) + .filter(t => isFillable(t.format)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kSamplePointMethods) .combine('C', ['i32', 'u32'] as const) - .combine('C_value', [-1, 0, 1, 2, 3, 4] as const) - .combine('coords', generateCoordBoundaries(2)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('offset', generateOffsets(2)) + .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('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); + }) + .fn(async t => { + const { format, 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 }); + 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, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + offset, + component: true, + hashInputs: [format, samplePoints, C, A, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, component, arrayIndex, offset }) => { + return { + builtin: 'textureGather', + component, + componentType: C === 'i32' ? 'i' : 'u', + coordType: 'f', + coords, + 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_array_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegather') @@ -140,8 +342,9 @@ g.test('sampled_array_3d_coords') ` C: i32, u32 T: i32, u32, f32 +A: i32, u32 -fn textureGather(component: C, t: texture_cube_array, s: sampler, coords: vec3, array_index: C) -> vec4 +fn textureGather(component: C, t: texture_cube_array, s: sampler, coords: vec3, array_index: A) -> vec4 Parameters: * component: @@ -154,17 +357,79 @@ Parameters: * array_index: The 0-based texture array index ` ) - .paramsSubcasesOnly( - u => - u - .combine('T', ['f32-only', 'i32', 'u32'] as const) - .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 */ + .params(u => + u + .combine('format', kTestableColorFormats) + .filter(t => isFillable(t.format)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('C', ['i32', 'u32'] as const) + .combine('A', ['i32', 'u32'] as const) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + t.skipIfTextureViewDimensionNotSupported('cube-array'); + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, t.params.format); + }) + .fn(async t => { + const { format, C, A, samplePoints, addressMode, minFilter } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube-array'; + const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + mipLevelCount: 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[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + component: true, + textureBuiltin: 'textureGather', + arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, + hashInputs: [format, C, samplePoints, addressMode, minFilter], + }).map(({ coords, component, arrayIndex }) => { + return { + builtin: 'textureGather', + component, + componentType: C === 'i32' ? 'i' : 'u', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + coordType: 'f', + coords, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = appendComponentTypeForFormatToTextureType('texture_cube_array', format); + 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/#texturegather') @@ -185,13 +450,67 @@ 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('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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .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('offset', [false, true] as const) ) - .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, + offset, + hashInputs: [format, samplePoints, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, offset }) => { + return { + builtin: 'textureGather', + coordType: 'f', + coords, + offset, + }; + }); + const textureType = 'texture_depth_2d'; + 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_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegather') @@ -205,21 +524,79 @@ Parameters: * coords: The texture coordinates ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('coords', generateCoordBoundaries(3)) + .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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) ) - .unimplemented(); + .fn(async t => { + const { format, samplePoints, addressMode, minFilter } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube'; + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + const depthOrArrayLayers = 6; + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size: { width, height, depthOrArrayLayers }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + mipLevelCount: 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[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + textureBuiltin: 'textureGather', + hashInputs: [format, samplePoints, addressMode, minFilter], + }).map(({ coords, component }) => { + return { + builtin: 'textureGather', + coordType: 'f', + coords, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = 'texture_depth_cube'; + 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/#texturegather') .desc( ` -C: i32, u32 +A: i32, u32 -fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: C) -> vec4 -fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: C, offset: vec2) -> vec4 +fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: A) -> vec4 +fn textureGather(t: texture_depth_2d_array, s: sampler, coords: vec2, array_index: A, offset: vec2) -> vec4 Parameters: * t: The depth texture to read from @@ -234,23 +611,85 @@ 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('coords', generateCoordBoundaries(2)) - /* array_index not param'd as out-of-bounds is implementation specific */ - .combine('offset', generateOffsets(2)) + .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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .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('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice(t, t.params.minFilter, 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, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + offset, + hashInputs: [format, samplePoints, A, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, arrayIndex, offset }) => { + return { + builtin: 'textureGather', + coordType: 'f', + coords, + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + offset, + }; + }); + const textureType = 'texture_depth_2d_array'; + 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_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegather') .desc( ` -C: i32, u32 +A: i32, u32 -fn textureGather(t: texture_depth_cube_array, s: sampler, coords: vec3, array_index: C) -> vec4 +fn textureGather(t: texture_depth_cube_array, s: sampler, coords: vec3, array_index: A) -> vec4 Parameters: * t: The depth texture to read from @@ -259,12 +698,73 @@ Parameters: * array_index: The 0-based texture array index ` ) - .paramsSubcasesOnly( - u => - u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('C', ['i32', 'u32'] as const) - .combine('coords', generateCoordBoundaries(3)) - /* array_index not param'd as out-of-bounds is implementation specific */ + .params(u => + u + .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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('A', ['i32', 'u32'] as const) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported('cube-array'); + }) + .fn(async t => { + const { format, A, samplePoints, addressMode, minFilter } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube-array'; + const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + mipLevelCount: 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[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + textureBuiltin: 'textureGather', + arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, + hashInputs: [format, samplePoints, addressMode, minFilter], + }).map(({ coords, component, arrayIndex }) => { + return { + builtin: 'textureGather', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + coordType: 'f', + coords, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = '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/textureLoad.spec.ts b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts index 879817ec8ca3..534bf17c5c6e 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureLoad.spec.ts @@ -167,8 +167,7 @@ Parameters: .params(u => u .combine('format', kTestableColorFormats) - // MAINTENANCE_TODO: Update createTextureFromTexelViews to support stencil8 and remove this filter. - .filter(t => t.format !== 'stencil8' && !isCompressedFloatTextureFormat(t.format)) + .filter(t => !isCompressedFloatTextureFormat(t.format)) .beginSubcases() .combine('samplePoints', kSamplePointMethods) .combine('C', ['i32', 'u32'] as const) @@ -188,10 +187,7 @@ Parameters: const descriptor: GPUTextureDescriptor = { format, size, - usage: - GPUTextureUsage.COPY_DST | - GPUTextureUsage.TEXTURE_BINDING | - (canUseAsRenderTarget(format) ? GPUTextureUsage.RENDER_ATTACHMENT : 0), + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, mipLevelCount: maxMipLevelCount({ size }), }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); @@ -422,10 +418,7 @@ Parameters: const descriptor: GPUTextureDescriptor = { format, size, - usage: - GPUTextureUsage.COPY_DST | - GPUTextureUsage.TEXTURE_BINDING | - GPUTextureUsage.RENDER_ATTACHMENT, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, mipLevelCount: maxMipLevelCount({ size }), }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); 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 729563553260..5c6e99eb9665 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureSampleLevel.spec.ts @@ -5,15 +5,6 @@ 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'; @@ -284,7 +275,7 @@ Parameters: .fn(async t => { const { format, viewDimension, samplePoints, addressMode, minFilter, offset } = t.params; - const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + const [width, height] = chooseTextureSize({ minSize: 32, minBlocks: 2, format, viewDimension }); const depthOrArrayLayers = getDepthOrArrayLayersForViewDimension(viewDimension); const descriptor: GPUTextureDescriptor = { @@ -293,8 +284,7 @@ Parameters: ...(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, + mipLevelCount: 3, }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); const sampler: GPUSamplerDescriptor = { @@ -396,7 +386,7 @@ Parameters: const viewDimension: GPUTextureViewDimension = 'cube-array'; const size = chooseTextureSize({ - minSize: 8, + minSize: 32, minBlocks: 4, format, viewDimension, @@ -405,8 +395,7 @@ Parameters: format, size, usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, - // MAINTENANCE_TODO: Set this to 3. See above. - mipLevelCount: 1, + mipLevelCount: 3, }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); const sampler: GPUSamplerDescriptor = { @@ -423,8 +412,8 @@ Parameters: sampler, descriptor, mipLevel: { num: texture.mipLevelCount, type: 'f32' }, - arrayIndex: { num: texture.depthOrArrayLayers, type: A }, - hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter], + arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, + hashInputs: [format, viewDimension, A, samplePoints, addressMode, minFilter], }).map(({ coords, mipLevel, arrayIndex }) => { return { builtin: 'textureSampleLevel', @@ -456,7 +445,7 @@ g.test('depth_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') .desc( ` -C is i32 or u32 +L is i32 or u32 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 @@ -504,10 +493,7 @@ Parameters: format, size: { width, height }, mipLevelCount: 3, - usage: - GPUTextureUsage.COPY_DST | - GPUTextureUsage.TEXTURE_BINDING | - GPUTextureUsage.RENDER_ATTACHMENT, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); const sampler: GPUSamplerDescriptor = { @@ -553,7 +539,8 @@ g.test('depth_array_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') .desc( ` -C is i32 or u32 +A is i32 or u32 +L is i32 or u32 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 @@ -603,10 +590,7 @@ Parameters: format, size: { width, height }, mipLevelCount: 3, - usage: - GPUTextureUsage.COPY_DST | - GPUTextureUsage.TEXTURE_BINDING | - GPUTextureUsage.RENDER_ATTACHMENT, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, ...(t.isCompatibility && { textureBindingViewDimension: '2d-array' }), }; const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); @@ -656,7 +640,8 @@ g.test('depth_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturesamplelevel') .desc( ` -C is i32 or u32 +L is i32 or u32 +A is i32 or u32 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 @@ -704,7 +689,7 @@ Parameters: const { format, viewDimension, samplePoints, A, L, addressMode, minFilter } = t.params; const size = chooseTextureSize({ - minSize: 8, + minSize: 32, minBlocks: 4, format, viewDimension, @@ -712,10 +697,7 @@ Parameters: const descriptor: GPUTextureDescriptor = { format, size, - usage: - GPUTextureUsage.COPY_DST | - GPUTextureUsage.TEXTURE_BINDING | - GPUTextureUsage.RENDER_ATTACHMENT, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, mipLevelCount: 3, ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), }; @@ -733,8 +715,8 @@ Parameters: method: samplePoints, sampler, descriptor, - mipLevel: { num: texture.mipLevelCount, type: L }, - arrayIndex: A ? { num: texture.depthOrArrayLayers, type: A } : undefined, + mipLevel: { num: texture.mipLevelCount - 1, type: L }, + arrayIndex: A ? { num: texture.depthOrArrayLayers / 6, type: A } : undefined, hashInputs: [format, viewDimension, samplePoints, addressMode, minFilter], }).map(({ coords, mipLevel, arrayIndex }) => { return { 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 3f1d778b52b1..4da0c3dda5cf 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -1,10 +1,11 @@ -import { keysOf } from '../../../../../../common/util/data_tables.js'; import { assert, range, unreachable } from '../../../../../../common/util/util.js'; import { EncodableTextureFormat, isCompressedFloatTextureFormat, isCompressedTextureFormat, isDepthOrStencilTextureFormat, + isDepthTextureFormat, + isStencilTextureFormat, kEncodableTextureFormats, kTextureFormatInfo, } from '../../../../../format_info.js'; @@ -77,16 +78,52 @@ export function getTextureTypeForTextureViewDimension(viewDimension: GPUTextureV } } +const is32Float = (format: GPUTextureFormat) => + format === 'r32float' || format === 'rg32float' || format === 'rgba32float'; + +/** + * Skips a subcase if the filter === 'linear' and the format is type + * 'unfilterable-float' and we cannot enable filtering. + */ +export function skipIfNeedsFilteringAndIsUnfilterableOrSelectDevice( + t: GPUTestSubcaseBatchState, + filter: GPUFilterMode, + format: GPUTextureFormat +) { + const features = new Set(); + features.add(kTextureFormatInfo[format].feature); + + if (filter === 'linear') { + t.skipIf(isDepthTextureFormat(format), 'depth texture are unfilterable'); + + const type = kTextureFormatInfo[format].color?.type; + if (type === 'unfilterable-float') { + assert(is32Float(format)); + features.add('float32-filterable'); + } + } + + if (features.size > 0) { + t.selectDeviceOrSkipTestCase(Array.from(features)); + } +} + +/** + * Returns if a texture format can be filled with random data. + */ +export function isFillable(format: GPUTextureFormat) { + // 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. + return !isCompressedTextureFormat(format) || !format.endsWith('float'); +} + /** * 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; + return canPotentiallyFilter && isFillable(format); } /** @@ -491,6 +528,7 @@ export type Dimensionality = vec1 | vec2 | vec3; type TextureCallArgKeys = keyof TextureCallArgs; const kTextureCallArgNames: readonly TextureCallArgKeys[] = [ + 'component', 'coords', 'arrayIndex', 'sampleIndex', @@ -501,6 +539,7 @@ const kTextureCallArgNames: readonly TextureCallArgKeys[] = [ ] as const; export interface TextureCallArgs { + component?: number; coords?: T; mipLevel?: number; arrayIndex?: number; @@ -510,12 +549,20 @@ export interface TextureCallArgs { offset?: T; } +export type TextureBuiltin = + | 'textureGather' + | 'textureLoad' + | 'textureSample' + | 'textureSampleBaseClampToEdge' + | 'textureSampleLevel'; + export interface TextureCall extends TextureCallArgs { - builtin: 'textureLoad' | 'textureSample' | 'textureSampleBaseClampToEdge' | 'textureSampleLevel'; + builtin: TextureBuiltin; coordType: 'f' | 'i' | 'u'; levelType?: 'i' | 'u' | 'f'; arrayIndexType?: 'i' | 'u'; sampleIndexType?: 'i' | 'u'; + componentType?: 'i' | 'u'; } const s_u32 = new Uint32Array(1); @@ -546,6 +593,9 @@ function getCallArgType( switch (argName) { case 'coords': return call.coordType; + case 'component': + assert(call.componentType !== undefined); + return call.componentType; case 'mipLevel': assert(call.levelType !== undefined); return call.levelType; @@ -718,6 +768,7 @@ export function softwareTextureReadMipLevel( }; switch (call.builtin) { + case 'textureGather': case 'textureSample': case 'textureSampleBaseClampToEdge': case 'textureSampleLevel': { @@ -748,7 +799,7 @@ export function softwareTextureReadMipLevel( const samples: { at: number[]; weight: number }[] = []; - const filter = sampler?.minFilter ?? 'nearest'; + const filter = call.builtin === 'textureGather' ? 'linear' : sampler?.minFilter ?? 'nearest'; switch (filter) { case 'linear': { // 'p0' is the lower texel for 'at' @@ -767,10 +818,11 @@ export function softwareTextureReadMipLevel( samples.push({ at: p1, weight: p1W[0] }); break; case 2: { - samples.push({ at: p0, weight: p0W[0] * p0W[1] }); - samples.push({ at: [p1[0], p0[1]], weight: p1W[0] * p0W[1] }); + // Note: These are ordered to match textureGather samples.push({ at: [p0[0], p1[1]], weight: p0W[0] * p1W[1] }); samples.push({ at: p1, weight: p1W[0] * p1W[1] }); + samples.push({ at: [p1[0], p0[1]], weight: p1W[0] * p0W[1] }); + samples.push({ at: p0, weight: p0W[0] * p0W[1] }); break; } case 3: { @@ -780,10 +832,11 @@ export function softwareTextureReadMipLevel( // the slice they'll be wrapped by wrapFaceCoordToCubeFaceAtEdgeBoundaries // below. if (isCube) { - samples.push({ at: p0, weight: p0W[0] * p0W[1] }); - samples.push({ at: [p1[0], p0[1], p0[2]], weight: p1W[0] * p0W[1] }); + // Note: These are ordered to match textureGather samples.push({ at: [p0[0], p1[1], p0[2]], weight: p0W[0] * p1W[1] }); samples.push({ at: p1, weight: p1W[0] * p1W[1] }); + samples.push({ at: [p1[0], p0[1], p0[2]], weight: p1W[0] * p0W[1] }); + samples.push({ at: p0, weight: p0W[0] * p0W[1] }); const ndx = getUnusedCubeCornerSampleIndex(textureSize[0], coords as vec3); if (ndx >= 0) { // # Issues with corners of cubemaps @@ -817,7 +870,16 @@ export function softwareTextureReadMipLevel( // I'm not sure what "average the values of the three available samples" // means. To me that would be (a+b+c)/3 or in other words, set all the // weights to 0.33333 but that's not what the M1 is doing. - unreachable('corners of cubemaps are not testable'); + // + // We could check that, given the 3 texels at the corner, if all 3 texels + // are the same value then the result must be the same value. Otherwise, + // the result must be between the 3 values. For now, the code that + // chooses test coordinates avoids corners. This has the restriction + // that the smallest mip level be at least 4x4 so there are some non + // corners to choose from. + unreachable( + `corners of cubemaps are not testable:\n ${describeTextureCall(call)}` + ); } } else { const p = [p0, p1]; @@ -847,6 +909,24 @@ export function softwareTextureReadMipLevel( unreachable(); } + if (call.builtin === 'textureGather') { + const componentNdx = call.component ?? 0; + assert(componentNdx >= 0 && componentNdx < 4); + assert(samples.length === 4); + const component = kRGBAComponents[componentNdx]; + const out: PerTexelComponent = {}; + samples.forEach((sample, i) => { + const c = isCube + ? wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize[0], sample.at as vec3) + : applyAddressModesToCoords(addressMode, textureSize, sample.at); + const v = load(c); + const rgba = convertPerTexelComponentToResultFormat(v, format); + out[kRGBAComponents[i]] = rgba[component]; + }); + + return out; + } + const out: PerTexelComponent = {}; const ss = []; for (const sample of samples) { @@ -1124,6 +1204,50 @@ function texelsApproximatelyEqual( return true; } +// If it's `textureGather` then we need to convert all values to one component. +// In other words, imagine the format is rg11b10ufloat. If it was +// `textureSample` we'd have `r11, g11, b10, a=1` but for `textureGather` +// +// component = 0 => `r11, r11, r11, r11` +// component = 1 => `g11, g11, g11, g11` +// component = 2 => `b10, b10, b10, b10` +// +// etc..., each from a different texel +// +// The Texel utils don't handle this. So if `component = 2` we take each value, +// copy it to the `B` component, run it through the texel utils so it returns +// the correct ULP for a 10bit float (not an 11 bit float). Then copy it back to +// the channel it came from. +function getULPFromZeroForComponents( + rgba: PerTexelComponent, + format: EncodableTextureFormat, + builtin: TextureBuiltin, + componentNdx?: number +): PerTexelComponent { + const rep = kTexelRepresentationInfo[format]; + if (builtin === 'textureGather') { + const out: PerTexelComponent = {}; + const component = kRGBAComponents[componentNdx ?? 0]; + const temp: PerTexelComponent = { R: 0, G: 0, B: 0, A: 1 }; + for (const comp of kRGBAComponents) { + temp[component] = rgba[comp]; + const texel = convertResultFormatToTexelViewFormat(temp, format); + const ulp = convertPerTexelComponentToResultFormat( + rep.bitsToULPFromZero(rep.numberToBits(texel)), + format + ); + out[comp] = ulp[component]; + } + return out; + } else { + const texel = convertResultFormatToTexelViewFormat(rgba, format); + return convertPerTexelComponentToResultFormat( + rep.bitsToULPFromZero(rep.numberToBits(texel)), + format + ); + } +} + /** * Checks the result of each call matches the expected result. */ @@ -1137,7 +1261,6 @@ export async function checkCallResults( ) { const errs: string[] = []; const format = texture.texels[0].format; - const rep = kTexelRepresentationInfo[format]; const size = reifyExtent3D(texture.descriptor.size); const maxFractionalDiff = sampler?.minFilter === 'linear' || @@ -1146,11 +1269,21 @@ export async function checkCallResults( ? getMaxFractionalDiffForTextureFormat(texture.descriptor.format) : 0; - for (let callIdx = 0; callIdx < calls.length && errs.length === 0; callIdx++) { + for (let callIdx = 0; callIdx < calls.length /*&& errs.length === 0*/; callIdx++) { const call = calls[callIdx]; const gotRGBA = results[callIdx]; const expectRGBA = softwareTextureReadLevel(t, call, texture, sampler, call.mipLevel ?? 0); + // 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. + if ( + isDepthOrStencilTextureFormat(format) && + call.builtin === 'textureGather' && + call.component! > 0 + ) { + continue; + } + if (texelsApproximatelyEqual(gotRGBA, expectRGBA, format, maxFractionalDiff)) { continue; } @@ -1159,65 +1292,85 @@ export async function checkCallResults( continue; } - 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) { - const g = got[component]!; - const e = expect[component]!; + const gULP = getULPFromZeroForComponents(gotRGBA, format, call.builtin, call.component); + const eULP = getULPFromZeroForComponents(expectRGBA, format, call.builtin, call.component); + + // from the spec: https://gpuweb.github.io/gpuweb/#reading-depth-stencil + // depth and stencil values are D, ?, ?, ? + const rgbaComponentsToCheck = + call.builtin === 'textureGather' || !isDepthOrStencilTextureFormat(format) + ? kRGBAComponents + : kRComponent; + + let bad = false; + const diffs = rgbaComponentsToCheck.map(component => { + const g = gotRGBA[component]!; + const e = expectRGBA[component]!; const absDiff = Math.abs(g - e); const ulpDiff = Math.abs(gULP[component]! - eULP[component]!); - const relDiff = absDiff / Math.max(Math.abs(g), Math.abs(e)); + assert(!Number.isNaN(ulpDiff)); + const maxAbs = Math.max(Math.abs(g), Math.abs(e)); + const relDiff = maxAbs > 0 ? absDiff / maxAbs : 0; if (ulpDiff > 3 && absDiff > maxFractionalDiff) { - const desc = describeTextureCall(call); - errs.push(`component was not as expected: + bad = true; + } + return { absDiff, relDiff, ulpDiff }; + }); + + const fix5 = (n: number) => n.toFixed(5); + const fix5v = (arr: number[]) => arr.map(v => fix5(v)).join(', '); + const rgbaToArray = (p: PerTexelComponent): number[] => + rgbaComponentsToCheck.map(component => p[component]!); + + if (bad) { + const desc = describeTextureCall(call); + errs.push(`result was not as expected: size: [${size.width}, ${size.height}, ${size.depthOrArrayLayers}] mipCount: ${texture.descriptor.mipLevelCount ?? 1} call: ${desc} // #${callIdx} - component: ${component} - got: ${g} - expected: ${e} - abs diff: ${absDiff.toFixed(4)} - rel diff: ${(relDiff * 100).toFixed(2)}% - ulp diff: ${ulpDiff} + got: ${fix5v(rgbaToArray(gotRGBA))} + expected: ${fix5v(rgbaToArray(expectRGBA))} + max diff: ${maxFractionalDiff} + abs diffs: ${fix5v(diffs.map(({ absDiff }) => absDiff))} + rel diffs: ${diffs.map(({ relDiff }) => `${(relDiff * 100).toFixed(2)}%`).join(', ')} + ulp diffs: ${diffs.map(({ ulpDiff }) => ulpDiff).join(', ')} `); - if (sampler) { - const expectedSamplePoints = [ - 'expected:', - ...(await identifySamplePoints(texture, (texels: TexelView[]) => { - return Promise.resolve( - softwareTextureReadLevel( - t, - call, - { - texels, - descriptor: texture.descriptor, - viewDescriptor: texture.viewDescriptor, - }, - sampler, - call.mipLevel ?? 0 - ) - ); - })), - ]; - const gotSamplePoints = [ - 'got:', - ...(await identifySamplePoints(texture, async (texels: TexelView[]) => { - const gpuTexture = createTextureFromTexelViews(t, texels, texture.descriptor); - const result = ( - await doTextureCalls(t, gpuTexture, texture.viewDescriptor, textureType, sampler, [ - call, - ]) - )[0]; - gpuTexture.destroy(); - return result; - })), - ]; - errs.push(' sample points:'); - errs.push(layoutTwoColumns(expectedSamplePoints, gotSamplePoints).join('\n')); - errs.push('', ''); - } + + if (sampler) { + const expectedSamplePoints = [ + 'expected:', + ...(await identifySamplePoints(texture, call, (texels: TexelView[]) => { + return Promise.resolve( + softwareTextureReadLevel( + t, + call, + { + texels, + descriptor: texture.descriptor, + viewDescriptor: texture.viewDescriptor, + }, + sampler, + call.mipLevel ?? 0 + ) + ); + })), + ]; + const gotSamplePoints = [ + 'got:', + ...(await identifySamplePoints(texture, call, async (texels: TexelView[]) => { + const gpuTexture = createTextureFromTexelViewsLocal(t, texels, texture.descriptor); + const result = ( + await doTextureCalls(t, gpuTexture, texture.viewDescriptor, textureType, sampler, [ + call, + ]) + )[0]; + gpuTexture.destroy(); + return result; + })), + ]; + errs.push(' sample points:'); + errs.push(layoutTwoColumns(expectedSamplePoints, gotSamplePoints).join('\n')); + errs.push('', ''); } } } @@ -1797,6 +1950,19 @@ export async function readTextureToTexelViews( return texelViews; } +function createTextureFromTexelViewsLocal( + t: GPUTest, + texelViews: TexelView[], + desc: Omit +): GPUTexture { + const modifiedDescriptor = { ...desc }; + // If it's a depth or stencil texture we need to render to it to fill it with data. + if (isDepthOrStencilTextureFormat(texelViews[0].format)) { + modifiedDescriptor.usage = desc.usage | GPUTextureUsage.RENDER_ATTACHMENT; + } + return createTextureFromTexelViews(t, texelViews, modifiedDescriptor); +} + /** * Fills a texture with random data and returns that data as * an array of TexelView. @@ -1825,14 +1991,14 @@ export async function createTextureWithRandomDataAndGetTexels( return { texture, texels }; } else { const texels = createRandomTexelViewMipmap(descriptor); - const texture = createTextureFromTexelViews(t, texels, descriptor); + const texture = createTextureFromTexelViewsLocal(t, texels, descriptor); return { texture, texels }; } } function valueIfAllComponentsAreEqual( c: PerTexelComponent, - componentOrder: TexelComponent[] + componentOrder: readonly TexelComponent[] ) { const s = new Set(componentOrder.map(component => c[component]!)); return s.size === 1 ? s.values().next().value : undefined; @@ -1927,8 +2093,9 @@ const kFaceNames = ['+x', '-x', '+y', '-y', '+z', '-z'] as const; * a: at: [7, 1], weights: [R: 0.75000] * b: at: [7, 2], weights: [R: 0.25000] */ -async function identifySamplePoints( +async function identifySamplePoints( texture: Texture, + call: TextureCall, run: (texels: TexelView[]) => Promise> ) { const info = texture.descriptor; @@ -1968,6 +2135,10 @@ async function identifySamplePoints( ) as EncodableTextureFormat; const rep = kTexelRepresentationInfo[format]; + const components = call.builtin === 'textureGather' ? kRGBAComponents : rep.componentOrder; + const convertResultAsAppropriate = + call.builtin === 'textureGather' ? (v: T) => v : convertResultFormatToTexelViewFormat; + // Identify all the texels that are sampled, and their weights. const sampledTexelWeights = new Map>(); const unclassifiedStack = [new Set(range(numTexels, v => v))]; @@ -1986,7 +2157,7 @@ async function identifySamplePoints( } // See if any of the texels in setA were sampled. - const results = convertResultFormatToTexelViewFormat( + const results = convertResultAsAppropriate( await run( range(mipLevelCount, mipLevel => TexelView.fromTexelsAsColors( @@ -2012,7 +2183,7 @@ async function identifySamplePoints( ), format ); - if (rep.componentOrder.some(c => results[c] !== 0)) { + if (components.some(c => results[c] !== 0)) { // One or more texels of setA were sampled. if (setA.size === 1) { // We identified a specific texel was sampled. @@ -2074,13 +2245,15 @@ async function identifySamplePoints( for (let layer = 0; layer < depthOrArrayLayers; ++layer) { const layerEntries = level[layer]; - if (!layerEntries) { - continue; - } const orderedTexelIndices: number[] = []; lines.push(''); - lines.push(`layer: ${layer}${isCube ? ` (${kFaceNames[layer]})` : ''}`); + const unSampled = layerEntries ? '' : 'un-sampled'; + lines.push(`layer: ${layer}${isCube ? ` (${kFaceNames[layer]})` : ''} ${unSampled}`); + + if (!layerEntries) { + continue; + } { let line = ' '; @@ -2133,11 +2306,11 @@ async function identifySamplePoints( const weights = layerEntries.get(texelIdx)!; const y = Math.floor(texelIdx / texelsPerRow); const x = texelIdx % texelsPerRow; - const singleWeight = valueIfAllComponentsAreEqual(weights, rep.componentOrder); + const singleWeight = valueIfAllComponentsAreEqual(weights, components); const w = singleWeight !== undefined ? `weight: ${fix5(singleWeight)}` - : `weights: [${rep.componentOrder.map(c => `${c}: ${fix5(weights[c]!)}`).join(', ')}]`; + : `weights: [${components.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}`); }); @@ -2213,11 +2386,13 @@ export const kCubeSamplePointMethods = ['cube-edges', 'texel-centre', 'spiral'] export type CubeSamplePointMethods = (typeof kSamplePointMethods)[number]; type TextureBuiltinInputArgs = { + textureBuiltin?: TextureBuiltin; descriptor: GPUTextureDescriptor; sampler?: GPUSamplerDescriptor; mipLevel?: RangeDef; sampleIndex?: RangeDef; arrayIndex?: RangeDef; + component?: boolean; offset?: boolean; hashInputs: (number | string | boolean)[]; }; @@ -2237,7 +2412,14 @@ function generateTextureBuiltinInputsImpl( radius?: number; loops?: number; }) -): { coords: T; mipLevel: number; sampleIndex?: number; arrayIndex?: number; offset?: T }[] { +): { + coords: T; + mipLevel: number; + sampleIndex?: number; + arrayIndex?: number; + offset?: T; + component?: number; +}[] { const { method, descriptor } = args; const dimension = descriptor.dimension ?? '2d'; const mipLevelCount = descriptor.mipLevelCount ?? 1; @@ -2270,7 +2452,7 @@ function generateTextureBuiltinInputsImpl( typeof v === 'string' ? sumOfCharCodesOfString(v) : typeof v === 'boolean' ? (v ? 1 : 0) : v ); const makeRangeValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => { - const range = num + type === 'u32' ? 1 : 2; + const range = num + (type === 'u32' ? 1 : 2); const number = (hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1); return type === 'f32' ? number : Math.floor(number); @@ -2290,7 +2472,12 @@ function generateTextureBuiltinInputsImpl( // Linux, AMD Radeon Pro WX 3200: 256 // MacOS, M1 Mac: 256 const kSubdivisionsPerTexel = 4; - const nearest = !args.sampler || args.sampler.minFilter === 'nearest'; + const avoidEdgeCase = + !args.sampler || + args.sampler.minFilter === 'nearest' || + args.textureBuiltin === 'textureGather'; + const edgeRemainder = args.textureBuiltin === 'textureGather' ? kSubdivisionsPerTexel / 2 : 0; + const numComponents = isDepthOrStencilTextureFormat(descriptor.format) ? 1 : 4; return coords.map((c, i) => { const mipLevel = args.mipLevel ? quantizeMipLevel(makeRangeValue(args.mipLevel, i), args.sampler?.mipmapFilter ?? 'nearest') @@ -2302,9 +2489,10 @@ function generateTextureBuiltinInputsImpl( const coords = c.map((v, i) => { // Quantize to kSubdivisionsPerPixel const v1 = Math.floor(v * q[i]); - // If it's nearest and we're on the edge of a texel then move us off the edge - // since the edge could choose one texel or another in nearest mode - const v2 = nearest && v1 % kSubdivisionsPerTexel === 0 ? v1 + 1 : v1; + // If it's nearest or textureGather and we're on the edge of a texel then move us off the edge + // since the edge could choose one texel or another. + const isEdgeCase = v1 % kSubdivisionsPerTexel === edgeRemainder; + const v2 = isEdgeCase && avoidEdgeCase ? v1 + 1 : v1; // Convert back to texture coords return v2 / q[i]; }) as T; @@ -2317,12 +2505,24 @@ function generateTextureBuiltinInputsImpl( offset: args.offset ? (coords.map((_, j) => makeIntHashValueRepeatable(-8, 8, i, 3 + j)) as T) : undefined, + component: args.component ? makeIntHashValueRepeatable(0, numComponents, i, 4) : undefined, }; }); } +/** + * When mipmapFilter === 'nearest' we need to stay away from 0.5 + * because the GPU could decide to choose one mip or the other. + * + * Some example transition values, the value at which the GPU chooses + * mip level 1 over mip level 0: + * + * M1 Mac: 0.515381 + * Intel Mac: 0.49999 + * AMD Mac: 0.5 + */ const kMipEpsilon = 0.02; -function quantizeMipLevel(mipLevel: number, mipmapFilter: GPUFilterMode) { +function quantizeMipLevel(mipLevel: number, mipmapFilter: GPUMipmapFilterMode) { if (mipmapFilter === 'linear') { return mipLevel; } @@ -2432,135 +2632,35 @@ export function convertNormalized3DTexCoordToCubeCoord(uvLayer: vec3) { } /** + * Wrap a texel based face coord across cube faces + * * We have a face texture in texels coord where U/V choose a texel and W chooses the face. * If U/V are outside the size of the texture then, when normalized and converted * to a cube map coordinate, they'll end up pointing to a different face. * * addressMode is effectively ignored for cube * - * +-----------+ - * |0->u | - * |↓ | - * |v +y | - * | (2) | - * | | - * +-----------+-----------+-----------+-----------+ - * |0->u |0->u |0->u |0->u | - * |↓ |↓ |↓ |↓ | - * |v -x |v +z |v +x |v -z | - * | (1) | (4) | (0) | (5) | - * | | | | | - * +-----------+-----------+-----------+-----------+ - * |0->u | - * |↓ | - * |v -y | - * | (3) | - * | | - * +-----------+ + * By converting from a texel based coord to a normalized coord and then to a cube map coord, + * if the texel was outside of the face, the cube map coord will end up pointing to a different + * face. We then convert back cube coord -> normalized face coord -> texel based coord */ -const kFaceConversions = { - u: (textureSize: number, faceCoord: vec3) => faceCoord[0], - v: (textureSize: number, faceCoord: vec3) => faceCoord[1], - 'u+t': (textureSize: number, faceCoord: vec3) => faceCoord[0] + textureSize, - 'u-t': (textureSize: number, faceCoord: vec3) => faceCoord[0] - textureSize, - 'v+t': (textureSize: number, faceCoord: vec3) => faceCoord[1] + textureSize, - 'v-t': (textureSize: number, faceCoord: vec3) => faceCoord[1] - textureSize, - 't-v': (textureSize: number, faceCoord: vec3) => textureSize - faceCoord[1], - '1+u': (textureSize: number, faceCoord: vec3) => 1 + faceCoord[0], - '1+v': (textureSize: number, faceCoord: vec3) => 1 + faceCoord[1], - '-v-1': (textureSize: number, faceCoord: vec3) => -faceCoord[1] - 1, - 't-u-1': (textureSize: number, faceCoord: vec3) => textureSize - faceCoord[0] - 1, - 't-v-1': (textureSize: number, faceCoord: vec3) => textureSize - faceCoord[1] - 1, - '2t-u-1': (textureSize: number, faceCoord: vec3) => textureSize * 2 - faceCoord[0] - 1, - '2t-v-1': (textureSize: number, faceCoord: vec3) => textureSize * 2 - faceCoord[1] - 1, -} as const; -const kFaceConversionEnums = keysOf(kFaceConversions); -type FaceCoordConversion = (typeof kFaceConversionEnums)[number]; - -// For Each face -// face to go if u < 0 -// face to go if u >= textureSize -// face to go if v < 0 -// face to go if v >= textureSize -const kFaceToFaceRemap: { to: number; u: FaceCoordConversion; v: FaceCoordConversion }[][] = [ - // 0 - [ - /* -u */ { to: 4, u: 'u+t', v: 'v' }, - /* +u */ { to: 5, u: 'u-t', v: 'v' }, - /* -v */ { to: 2, u: 'v+t', v: 't-u-1' }, - /* +v */ { to: 3, u: '2t-v-1', v: 'u' }, - ], - // 1 - [ - /* -u */ { to: 5, u: 'u+t', v: 'v' }, - /* +u */ { to: 4, u: 'u-t', v: 'v' }, - /* -v */ { to: 2, u: '-v-1', v: 'u' }, // -1->0, -2->1 -3->2 - /* +v */ { to: 3, u: 't-v', v: 't-u-1' }, - ], - // 2 - [ - /* -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: '-v-1' }, - /* +v */ { to: 4, u: 'u', v: 'v-t' }, - ], - // 3 - [ - /* -u */ { to: 1, u: 't-v-1', v: 'u+t' }, - /* +u */ { to: 0, u: 'v', v: '2t-u-1' }, - /* -v */ { to: 4, u: 'u', v: 'v+t' }, - /* +v */ { to: 5, u: 't-u-1', v: '2t-v-1' }, - ], - // 4 - [ - /* -u */ { to: 1, u: 'u+t', v: 'v' }, - /* +u */ { to: 0, u: 'u-t', v: 'v' }, - /* -v */ { to: 2, u: 'u', v: 'v+t' }, - /* +v */ { to: 3, u: 'u', v: 'v-t' }, - ], - // 5 - [ - /* -u */ { to: 0, u: 'u+t', v: 'v' }, - /* +u */ { to: 1, u: 'u-t', v: 'v' }, - /* -v */ { to: 2, u: 't-u-1', v: '1+v' }, - /* +v */ { to: 3, u: 't-u-1', v: '2t-v-1' }, - ], -]; - -function getFaceWrapIndex(textureSize: number, faceCoord: vec3) { - if (faceCoord[0] < 0) { - return 0; - } - if (faceCoord[0] >= textureSize) { - return 1; - } - if (faceCoord[1] < 0) { - return 2; - } - if (faceCoord[1] >= textureSize) { - return 3; - } - return -1; -} - -function applyFaceWrap(textureSize: number, faceCoord: vec3): vec3 { - const ndx = getFaceWrapIndex(textureSize, faceCoord); - if (ndx < 0) { - return faceCoord; - } - const { to, u, v } = kFaceToFaceRemap[faceCoord[2]][ndx]; - return [ - kFaceConversions[u](textureSize, faceCoord), - kFaceConversions[v](textureSize, faceCoord), - to, +function wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize: number, faceCoord: vec3) { + // convert texel based face coord to normalized 2d-array coord + const nc0: vec3 = [ + (faceCoord[0] + 0.5) / textureSize, + (faceCoord[1] + 0.5) / textureSize, + (faceCoord[2] + 0.5) / 6, + ]; + const cc = convertNormalized3DTexCoordToCubeCoord(nc0); + const nc1 = convertCubeCoordToNormalized3DTextureCoord(cc); + // convert normalized 2d-array coord back texel based face coord + const fc = [ + Math.floor(nc1[0] * textureSize), + Math.floor(nc1[1] * textureSize), + Math.floor(nc1[2] * 6), ]; -} -function wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize: number, faceCoord: vec3) { - // If we're off both edges we need to wrap twice, once for each edge. - const faceCoord1 = applyFaceWrap(textureSize, faceCoord); - const faceCoord2 = applyFaceWrap(textureSize, faceCoord1); - return faceCoord2; + return fc; } function applyAddressModesToCoords( @@ -2607,6 +2707,7 @@ export function generateSamplePointsCube( mipLevel: number; arrayIndex?: number; offset?: undefined; + component?: number; }[] { const { method, descriptor } = args; const mipLevelCount = descriptor.mipLevelCount ?? 1; @@ -2647,20 +2748,38 @@ export function generateSamplePointsCube( /* prettier-ignore */ coords.push( // between edges - [-1.01, -1.02, 0], - [ 1.01, -1.02, 0], - [-1.01, 1.02, 0], - [ 1.01, 1.02, 0], - - [-1.01, 0, -1.02], - [ 1.01, 0, -1.02], - [-1.01, 0, 1.02], - [ 1.01, 0, 1.02], - - [-1.01, -1.02, 0], - [ 1.01, -1.02, 0], - [-1.01, 1.02, 0], - [ 1.01, 1.02, 0], + // +x + [ 1 , -1.01, 0 ], // wrap -y + [ 1 , +1.01, 0 ], // wrap +y + [ 1 , 0 , -1.01 ], // wrap -z + [ 1 , 0 , +1.01 ], // wrap +z + // -x + [ -1 , -1.01, 0 ], // wrap -y + [ -1 , +1.01, 0 ], // wrap +y + [ -1 , 0 , -1.01 ], // wrap -z + [ -1 , 0 , +1.01 ], // wrap +z + + // +y + [ -1.01, 1 , 0 ], // wrap -x + [ +1.01, 1 , 0 ], // wrap +x + [ 0 , 1 , -1.01 ], // wrap -z + [ 0 , 1 , +1.01 ], // wrap +z + // -y + [ -1.01, -1 , 0 ], // wrap -x + [ +1.01, -1 , 0 ], // wrap +x + [ 0 , -1 , -1.01 ], // wrap -z + [ 0 , -1 , +1.01 ], // wrap +z + + // +z + [ -1.01, 0 , 1 ], // wrap -x + [ +1.01, 0 , 1 ], // wrap +x + [ 0 , -1.01, 1 ], // wrap -y + [ 0 , +1.01, 1 ], // wrap +y + // -z + [ -1.01, 0 , -1 ], // wrap -x + [ +1.01, 0 , -1 ], // wrap +x + [ 0 , -1.01, -1 ], // wrap -y + [ 0 , +1.01, -1 ], // wrap +y // corners (see comment "Issues with corners of cubemaps") // for why these are commented out. @@ -2681,11 +2800,15 @@ export function generateSamplePointsCube( typeof v === 'string' ? sumOfCharCodesOfString(v) : typeof v === 'boolean' ? (v ? 1 : 0) : v ); const makeRangeValue = ({ num, type }: RangeDef, ...hashInputs: number[]) => { - const range = num + type === 'u32' ? 1 : 2; + const range = num + (type === 'u32' ? 1 : 2); const number = (hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range - (type === 'u32' ? 0 : 1); return type === 'f32' ? number : Math.floor(number); }; + const makeIntHashValue = (min: number, max: number, ...hashInputs: number[]) => { + const range = max - min; + return min + Math.floor((hashU32(..._hashInputs, ...hashInputs) / 0x1_0000_0000) * range); + }; // Samplers across devices use different methods to interpolate. // Quantizing the texture coordinates seems to hit coords that produce @@ -2695,12 +2818,105 @@ export function generateSamplePointsCube( // Win 11, NVidia 2070 Super: 16 // Linux, AMD Radeon Pro WX 3200: 256 // MacOS, M1 Mac: 256 + // + // Note: When doing `textureGather...` we can't use texel centers + // because which 4 pixels will be gathered jumps if we're slightly under + // or slightly over the center + // + // Similarly, if we're using 'nearest' filtering then we don't want texel + // edges for the same reason. + // + // Also note that for textureGather. The way it works for cube maps is to + // first convert from cube map coordinate to a 2D texture coordinate and + // a face. Then, choose 4 texels just like normal 2D texture coordinates. + // If one of the 4 texels is outside the current face, wrap it to the correct + // face. + // + // An issue this brings up though. Imagine a 2D texture with addressMode = 'repeat' + // + // 2d texture (same texture repeated to show 'repeat') + // ┌───┬───┬───┐ ┌───┬───┬───┐ + // │ │ │ │ │ │ │ │ + // ├───┼───┼───┤ ├───┼───┼───┤ + // │ │ │ a│ │c │ │ │ + // ├───┼───┼───┤ ├───┼───┼───┤ + // │ │ │ b│ │d │ │ │ + // └───┴───┴───┘ └───┴───┴───┘ + // + // Assume the texture coordinate is at the bottom right corner of a. + // Then textureGather will grab c, d, b, a (no idea why that order). + // but think of it as top-right, bottom-right, bottom-left, top-left. + // Similarly, if the texture coordinate is at the top left of d it + // will select the same 4 texels. + // + // But, in the case of a cubemap, each face is in different direction + // relative to the face next to it. + // + // +-----------+ + // |0->u | + // |↓ | + // |v +y | + // | (2) | + // | | + // +-----------+-----------+-----------+-----------+ + // |0->u |0->u |0->u |0->u | + // |↓ |↓ |↓ |↓ | + // |v -x |v +z |v +x |v -z | + // | (1) | (4) | (0) | (5) | + // | | | | | + // +-----------+-----------+-----------+-----------+ + // |0->u | + // |↓ | + // |v -y | + // | (3) | + // | | + // +-----------+ + // + // As an example, imagine going from the +y to the +x face. + // See diagram above, the right edge of the +y face wraps + // to the top edge of the +x face. + // + // +---+---+ + // | a|c | + // ┌───┬───┬───┐ ┌───┬───┬───┐ + // │ │ │ │ │ b│d │ │ + // ├───┼───┼───┤---+ ├───┼───┼───┤ + // │ │ │ a│ c | │ │ │ │ + // ├───┼───┼───┤---+ ├───┼───┼───┤ + // │ │ │ b│ d | │ │ │ │ + // └───┴───┴───┘---+ └───┴───┴───┘ + // +y face +x face + // + // If the texture coordinate is in the bottom right corner of a, + // the rectangle of texels we read are a,b,c,d and, if we the + // texture coordinate is in the top left corner of d we also + // read a,b,c,d according to the 2 diagrams above. + // + // But, notice that when reading from the POV of +y vs +x, + // which actual a,b,c,d texels are different. + // + // From the POV of face +x: a,b are in face +x and c,d are in face +y + // From the POV of face +y: a,c are in face +x and b,d are in face +y + // + // This is all the long way of saying that if we're on the edge of a cube + // face we could get drastically different results because the orientation + // of the rectangle of the 4 texels we use, rotates. So, we need to avoid + // any values too close to the edge just in case our math is different than + // the GPU's. + // const kSubdivisionsPerTexel = 4; - const nearest = !args.sampler || args.sampler.minFilter === 'nearest'; + const avoidEdgeCase = + !args.sampler || + args.sampler.minFilter === 'nearest' || + args.textureBuiltin === 'textureGather'; + const edgeRemainder = args.textureBuiltin === 'textureGather' ? kSubdivisionsPerTexel / 2 : 0; + return coords.map((c, i) => { - const mipLevel = args.mipLevel ? makeRangeValue(args.mipLevel, i) : 0; + 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('2d', size, clampedMipLevel); + const mipSize = virtualMipSize('2d', size, Math.ceil(clampedMipLevel)); const q = [ mipSize[0] * kSubdivisionsPerTexel, mipSize[0] * kSubdivisionsPerTexel, @@ -2720,17 +2936,20 @@ export function generateSamplePointsCube( const quantizedUVW = uvw.map((v, i) => { // Quantize to kSubdivisionsPerPixel const v1 = Math.floor(v * q[i]); - // If it's nearest and we're on the edge of a texel then move us off the edge - // since the edge could choose one texel or another in nearest mode - const v2 = nearest && v1 % kSubdivisionsPerTexel === 0 ? v1 + 1 : v1; - // Convert back to texture coords - return v2 / q[i]; + // If it's nearest or textureGather and we're on the edge of a texel then move us off the edge + // since the edge could choose one texel or another. + const isEdgeCase = v1 % kSubdivisionsPerTexel === edgeRemainder; + const v2 = isEdgeCase && avoidEdgeCase ? v1 + 1 : v1; + // Convert back to texture coords slightly off + return (v2 + 1 / 32) / q[i]; }) as vec3; + const coords = convertNormalized3DTexCoordToCubeCoord(quantizedUVW); return { coords, mipLevel, arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined, + component: args.component ? makeIntHashValue(0, 4, i, 4) : undefined, }; }); } @@ -2788,8 +3007,8 @@ function binKey(call: TextureCall): string { for (const name of kTextureCallArgNames) { const value = call[name]; if (value !== undefined) { - if (name === 'offset') { - // offset must be a constant expression + if (name === 'offset' || name === 'component') { + // offset and component must be constant expressions keys.push(`${name}: ${wgslExpr(value)}`); } else { keys.push(`${name}: ${wgslTypeFor(value, call.coordType)}`); @@ -2800,12 +3019,22 @@ function binKey(call: TextureCall): string { } function buildBinnedCalls(calls: TextureCall[]) { - const args: string[] = ['T']; // All texture builtins take the texture as the first argument + const args: string[] = []; const fields: string[] = []; const data: number[] = []; - const prototype = calls[0]; - if (prototype.builtin.startsWith('textureSample')) { + + if (prototype.builtin.startsWith('textureGather') && prototype['componentType']) { + args.push(`/* component */ ${wgslExpr(prototype['component']!)}`); + } + + // All texture builtins take a Texture + args.push('T'); + + if ( + prototype.builtin.startsWith('textureSample') || + prototype.builtin.startsWith('textureGather') + ) { // textureSample*() builtins take a sampler as the second argument args.push('S'); } @@ -2815,6 +3044,8 @@ function buildBinnedCalls(calls: TextureCall[]) { if (value !== undefined) { if (name === 'offset') { args.push(`/* offset */ ${wgslExpr(value)}`); + } else if (name === 'component') { + // was handled above } else { const type = name === 'mipLevel' @@ -2837,7 +3068,7 @@ function buildBinnedCalls(calls: TextureCall[]) { (prototype[name] === undefined) === (value === undefined), 'texture calls are not binned correctly' ); - if (value !== undefined && name !== 'offset') { + if (value !== undefined && name !== 'offset' && name !== 'component') { const type = getCallArgType(call, name); const bitcastToU32 = kBitCastFunctions[type]; if (value instanceof Array) { @@ -2877,13 +3108,17 @@ function binCalls(calls: TextureCall[]): number[][] } export function describeTextureCall(call: TextureCall): string { - const args: string[] = ['texture: T']; - if (call.builtin.startsWith('textureSample')) { + const args: string[] = []; + if (call.builtin.startsWith('textureGather') && call.componentType) { + args.push(`component: ${wgslExprFor(call.component!, call.componentType)}`); + } + args.push('texture: T'); + if (call.builtin.startsWith('textureSample') || call.builtin.startsWith('textureGather')) { args.push('sampler: S'); } for (const name of kTextureCallArgNames) { const value = call[name]; - if (value !== undefined) { + if (value !== undefined && name !== 'component') { if (name === 'coords') { args.push(`${name}: ${wgslExprFor(value, call.coordType)}`); } else if (name === 'mipLevel') { @@ -2922,6 +3157,21 @@ export async function doTextureCalls( sampler: GPUSamplerDescriptor | undefined, calls: TextureCall[] ) { + const { + format, + dimension, + depthOrArrayLayers, + sampleCount, + }: { + format: GPUTextureFormat; + dimension: GPUTextureDimension; + depthOrArrayLayers: number; + sampleCount: number; + } = + gpuTexture instanceof GPUExternalTexture + ? { format: 'rgba8unorm', dimension: '2d', depthOrArrayLayers: 1, sampleCount: 1 } + : gpuTexture; + let structs = ''; let body = ''; let dataFields = ''; @@ -2955,11 +3205,13 @@ export async function doTextureCalls( t.device.queue.writeBuffer(dataBuffer, 0, new Uint32Array(data)); const { resultType, resultFormat, componentType } = - gpuTexture instanceof GPUExternalTexture + calls[0].builtin === 'textureGather' + ? getTextureFormatTypeInfo(format) + : gpuTexture instanceof GPUExternalTexture ? ({ resultType: 'vec4f', resultFormat: 'rgba32float', componentType: 'f32' } as const) : textureType.includes('depth') ? ({ resultType: 'f32', resultFormat: 'rgba32float', componentType: 'f32' } as const) - : getTextureFormatTypeInfo(gpuTexture.format); + : getTextureFormatTypeInfo(format); const returnType = `vec4<${componentType}>`; const rtWidth = 256; @@ -3001,13 +3253,98 @@ ${body} const pipelines = s_deviceToPipelines.get(t.device) ?? new Map(); s_deviceToPipelines.set(t.device, pipelines); - const id = `${renderTarget.format}:${code}`; + // unfilterable-float textures can only be used with manually created bindGroupLayouts + // since the default 'auto' layout requires filterable textures/samplers. + // So, if we don't need filtering, don't request a filtering sampler. If we require + // filtering then check if the format is 32float format and if float32-filterable + // is enabled. + const info = kTextureFormatInfo[format ?? 'rgba8unorm']; + const isFiltering = + !!sampler && + (sampler.minFilter === 'linear' || + sampler.magFilter === 'linear' || + sampler.mipmapFilter === 'linear'); + let sampleType: GPUTextureSampleType = textureType.startsWith('texture_depth') + ? 'depth' + : isDepthTextureFormat(format) + ? 'unfilterable-float' + : isStencilTextureFormat(format) + ? 'uint' + : info.color?.type ?? 'float'; + if (isFiltering && sampleType === 'unfilterable-float') { + assert(is32Float(format)); + assert(t.device.features.has('float32-filterable')); + sampleType = 'float'; + } + if (sampleCount > 1 && sampleType === 'float') { + sampleType = 'unfilterable-float'; + } + + const entries: GPUBindGroupLayoutEntry[] = [ + { + binding: 2, + visibility: GPUShaderStage.FRAGMENT, + buffer: { + type: 'read-only-storage', + }, + }, + ]; + + const viewDimension = effectiveViewDimensionForDimension( + viewDescriptor.dimension, + dimension, + depthOrArrayLayers + ); + + if (textureType.includes('storage')) { + entries.push({ + binding: 0, + visibility: GPUShaderStage.FRAGMENT, + storageTexture: { + access: 'read-only', + viewDimension, + format, + }, + }); + } else if (gpuTexture instanceof GPUExternalTexture) { + entries.push({ + binding: 0, + visibility: GPUShaderStage.FRAGMENT, + externalTexture: {}, + }); + } else { + entries.push({ + binding: 0, + visibility: GPUShaderStage.FRAGMENT, + texture: { + sampleType, + viewDimension, + multisampled: sampleCount > 1, + }, + }); + } + + if (sampler) { + entries.push({ + binding: 1, + visibility: GPUShaderStage.FRAGMENT, + sampler: { + type: isFiltering ? 'filtering' : 'non-filtering', + }, + }); + } + + const id = `${renderTarget.format}:${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], + }); pipeline = await t.device.createRenderPipelineAsync({ - layout: 'auto', + layout, vertex: { module: shaderModule }, fragment: { module: shaderModule, diff --git a/src/webgpu/util/texture.ts b/src/webgpu/util/texture.ts index badce71baa34..20e99fdfad4d 100644 --- a/src/webgpu/util/texture.ts +++ b/src/webgpu/util/texture.ts @@ -17,6 +17,7 @@ const kLoadValueFromStorageInfo: Partial<{ texelType: string; unpackWGSL: string; useFragDepth?: boolean; + discardWithStencil?: boolean; }; }> = { r8unorm: { @@ -233,17 +234,27 @@ const kLoadValueFromStorageInfo: Partial<{ `, useFragDepth: true, }, + stencil8: { + storageType: 'u32', + texelType: 'vec4u', + unpackWGSL: ` + return vec4u(unpack4xU8(src[byteOffset / 4])[byteOffset % 4], 123, 123, 123) + `, + discardWithStencil: true, + }, }; function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { const info = kLoadValueFromStorageInfo[format]; assert(!!info); - const { storageType, texelType, unpackWGSL, useFragDepth } = info; + const { storageType, texelType, unpackWGSL, useFragDepth, discardWithStencil } = info; const [depthDecl, depthCode] = useFragDepth ? ['@builtin(frag_depth) d: f32,', 'fs.d = fs.v[0];'] : ['', '']; + const stencilCode = discardWithStencil ? 'if ((fs.v.r & vin.stencilMask) == 0) { discard; }' : ''; + return ` struct Uniforms { numTexelRows: u32, @@ -255,9 +266,10 @@ function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { struct VSOutput { @builtin(position) pos: vec4f, @location(0) @interpolate(flat, either) sampleIndex: u32, + @location(1) @interpolate(flat, either) stencilMask: u32, }; - @vertex fn vs(@builtin(vertex_index) vNdx: u32) -> VSOutput { + @vertex fn vs(@builtin(vertex_index) vNdx: u32, @builtin(instance_index) iNdx: u32) -> VSOutput { let points = array( vec2f(0, 0), vec2f(1, 0), vec2f(0, 1), vec2f(1, 1), ); @@ -266,7 +278,10 @@ function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { let rowOffset = f32(sampleRow) / numSampleRows; let rowMult = 1.0 / numSampleRows; let p = (points[vNdx % 4] * vec2f(1, rowMult) + vec2f(0, rowOffset)) * 2.0 - 1.0; - return VSOutput(vec4f(p, 0, 1), uni.sampleCount - sampleRow % uni.sampleCount - 1); + return VSOutput( + vec4f(p, 0, 1), + uni.sampleCount - sampleRow % uni.sampleCount - 1, + 1u << iNdx); } @group(0) @binding(0) var uni: Uniforms; @@ -289,6 +304,7 @@ function getCopyBufferToTextureViaRenderCode(format: GPUTextureFormat) { var fs: FSOutput; fs.v = unpack(byteOffset); ${depthCode} + ${stencilCode} return fs; } `; @@ -312,114 +328,158 @@ function copyBufferToTextureViaRender( const msInfo = kLoadValueFromStorageInfo[format]; assert(!!msInfo); - const { useFragDepth } = msInfo; + const { useFragDepth, discardWithStencil } = msInfo; const { device } = t; - const code = getCopyBufferToTextureViaRenderCode(format); - 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 } }), + const numBlits = discardWithStencil ? 8 : 1; + for (let blitCount = 0; blitCount < numBlits; ++blitCount) { + const code = getCopyBufferToTextureViaRenderCode(format); + const stencilWriteMask = 1 << blitCount; + const id = JSON.stringify({ + format, + useFragDepth, + stencilWriteMask, + discardWithStencil, + sampleCount, + code, }); - pipelines.set(id, pipeline); - } + 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({ + label: `blitCopyFor-${format}`, + layout: 'auto', + vertex: { module }, + ...(discardWithStencil + ? { + fragment: { + module, + targets: [], + }, + depthStencil: { + depthWriteEnabled: false, + depthCompare: 'always', + format, + stencilWriteMask, + stencilFront: { + passOp: 'replace', + }, + }, + } + : 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([ - copySize.height, // numTexelRows: u32, - source.bytesPerRow!, // bytesPerRow: u32, - info.bytesPerBlock!, // bytesPerSample: u32, - dest.texture.sampleCount, // sampleCount: u32, - ]); - const uniformBuffer = t.makeBufferWithContents( - uniforms, - GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM - ); - const storageBuffer = t.createBufferTracked({ - size: source.buffer.size, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, - }); - encoder.copyBufferToBuffer(source.buffer, 0, storageBuffer, 0, storageBuffer.size); - const baseMipLevel = dest.mipLevel; - for (let l = 0; l < copySize.depthOrArrayLayers; ++l) { - const baseArrayLayer = origin.z + l; - const mipLevelCount = 1; - const arrayLayerCount = 1; - const pass = encoder.beginRenderPass( - useFragDepth - ? { - colorAttachments: [], - depthStencilAttachment: { - view: dest.texture.createView({ - baseMipLevel, - baseArrayLayer, - mipLevelCount, - arrayLayerCount, - }), - depthClearValue: 0, - depthLoadOp: 'clear', - depthStoreOp: 'store', - }, - } - : { - colorAttachments: [ - { + const info = kTextureFormatInfo[format]; + const uniforms = new Uint32Array([ + copySize.height, // numTexelRows: u32, + source.bytesPerRow!, // bytesPerRow: u32, + info.bytesPerBlock!, // bytesPerSample: u32, + dest.texture.sampleCount, // sampleCount: u32, + ]); + const uniformBuffer = t.makeBufferWithContents( + uniforms, + GPUBufferUsage.COPY_DST | GPUBufferUsage.UNIFORM + ); + const storageBuffer = t.createBufferTracked({ + size: source.buffer.size, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, + }); + encoder.copyBufferToBuffer(source.buffer, 0, storageBuffer, 0, storageBuffer.size); + const baseMipLevel = dest.mipLevel; + for (let l = 0; l < copySize.depthOrArrayLayers; ++l) { + const baseArrayLayer = origin.z + l; + const mipLevelCount = 1; + const arrayLayerCount = 1; + const pass = encoder.beginRenderPass( + discardWithStencil + ? { + colorAttachments: [], + depthStencilAttachment: { view: dest.texture.createView({ baseMipLevel, baseArrayLayer, mipLevelCount, arrayLayerCount, }), - loadOp: 'clear', - storeOp: 'store', + stencilClearValue: 0, + stencilLoadOp: 'load', + stencilStoreOp: 'store', }, - ], - } - ); - pass.setViewport(origin.x, origin.y, copySize.width, copySize.height, 0, 1); - pass.setPipeline(pipeline); + } + : useFragDepth + ? { + colorAttachments: [], + depthStencilAttachment: { + view: dest.texture.createView({ + baseMipLevel, + baseArrayLayer, + mipLevelCount, + arrayLayerCount, + }), + depthClearValue: 0, + depthLoadOp: 'clear', + depthStoreOp: 'store', + }, + } + : { + colorAttachments: [ + { + view: dest.texture.createView({ + baseMipLevel, + baseArrayLayer, + mipLevelCount, + arrayLayerCount, + }), + loadOp: 'clear', + storeOp: 'store', + }, + ], + } + ); + pass.setViewport(origin.x, origin.y, copySize.width, copySize.height, 0, 1); + pass.setPipeline(pipeline); - const offset = - (source.offset ?? 0) + (source.bytesPerRow ?? 0) * (source.rowsPerImage ?? 0) * l; - const bindGroup = device.createBindGroup({ - layout: pipeline.getBindGroupLayout(0), - entries: [ - { binding: 0, resource: { buffer: uniformBuffer } }, - { binding: 1, resource: { buffer: storageBuffer, offset } }, - ], - }); + const offset = + (source.offset ?? 0) + (source.bytesPerRow ?? 0) * (source.rowsPerImage ?? 0) * l; + const bindGroup = device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: uniformBuffer } }, + { binding: 1, resource: { buffer: storageBuffer, offset } }, + ], + }); - pass.setBindGroup(0, bindGroup); - pass.draw(4 * copySize.height * dest.texture.sampleCount); - pass.end(); + pass.setBindGroup(0, bindGroup); + pass.setStencilReference(0xff); + pass.draw(4 * copySize.height * dest.texture.sampleCount, 1, 0, blitCount); + pass.end(); + } } } From 479198bf0fce0cd108396c88906829c0b5202692 Mon Sep 17 00:00:00 2001 From: Greggman Date: Fri, 23 Aug 2024 07:40:16 +0900 Subject: [PATCH 17/99] WGSL textureGatherCompare tests (#3914) --- src/webgpu/listing_meta.json | 4 +- .../call/builtin/textureGather.spec.ts | 2 +- .../call/builtin/textureGatherCompare.spec.ts | 332 ++++++++++++++++-- .../expression/call/builtin/texture_utils.ts | 111 ++++-- 4 files changed, 379 insertions(+), 70 deletions(-) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 2374c4cff7e1..4ed6e82ee734 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1554,8 +1554,8 @@ "webgpu:shader,execution,expression,call,builtin,textureGather:sampled_array_3d_coords:*": { "subcaseMS": 60.700 }, "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:array_2d_coords:*": { "subcaseMS": 291.301 }, "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:array_3d_coords:*": { "subcaseMS": 191.101 }, - "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_array_2d_coords:*": { "subcaseMS": 57.600 }, - "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_array_3d_coords:*": { "subcaseMS": 10.101 }, + "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_2d_coords:*": { "subcaseMS": 57.600 }, + "webgpu:shader,execution,expression,call,builtin,textureGatherCompare:sampled_3d_coords:*": { "subcaseMS": 10.101 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:arrayed:*": { "subcaseMS": 30.501 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:depth:*": { "subcaseMS": 3.200 }, "webgpu:shader,execution,expression,call,builtin,textureLoad:external:*": { "subcaseMS": 1.401 }, 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 f0f955b87dc8..9c062f7465f5 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGather.spec.ts @@ -744,7 +744,7 @@ Parameters: textureBuiltin: 'textureGather', arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, hashInputs: [format, samplePoints, addressMode, minFilter], - }).map(({ coords, component, arrayIndex }) => { + }).map(({ coords, arrayIndex }) => { return { builtin: 'textureGather', arrayIndex, 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 c743883ce849..6758e75f0018 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/textureGatherCompare.spec.ts @@ -17,20 +17,38 @@ A texture gather compare operation performs a depth comparison on four texels in `; import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; -import { GPUTest } from '../../../../../gpu_test.js'; +import { kCompareFunctions } from '../../../../../capability_info.js'; +import { + isDepthTextureFormat, + isEncodableTextureFormat, + kDepthStencilFormats, +} from '../../../../../format_info.js'; -import { generateCoordBoundaries, generateOffsets } from './utils.js'; +import { + checkCallResults, + chooseTextureSize, + createTextureWithRandomDataAndGetTexels, + doTextureCalls, + generateSamplePointsCube, + generateTextureBuiltinInputs2D, + kCubeSamplePointMethods, + kSamplePointMethods, + TextureCall, + vec2, + vec3, + WGSLTextureSampleTest, +} from './texture_utils.js'; -export const g = makeTestGroup(GPUTest); +export const g = makeTestGroup(WGSLTextureSampleTest); g.test('array_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegathercompare') .desc( ` -C: i32, u32 +A: i32, u32 -fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2, array_index: C, depth_ref: f32) -> vec4 -fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2, array_index: C, depth_ref: f32, offset: vec2) -> vec4 +fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2, array_index: A, depth_ref: f32) -> vec4 +fn textureGatherCompare(t: texture_depth_2d_array, s: sampler_comparison, coords: vec2, array_index: A, depth_ref: f32, offset: vec2) -> vec4 Parameters: * t: The depth texture to read from @@ -46,24 +64,87 @@ 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]) - .combine('coords', generateCoordBoundaries(2)) - .combine('depth_ref', [-1 /* smaller ref */, 0 /* equal ref */, 1 /* larger ref */] as const) - .combine('offset', generateOffsets(2)) + .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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .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('compare', kCompareFunctions) + .combine('offset', [false, true] as const) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureFormatNotSupported(t.params.format); + }) + .fn(async t => { + const { format, samplePoints, A, addressModeU, addressModeV, minFilter, compare, offset } = + t.params; + + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const depthOrArrayLayers = 4; + + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height, depthOrArrayLayers }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU, + addressModeV, + compare, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + arrayIndex: { num: texture.depthOrArrayLayers, type: A }, + depthRef: true, + offset, + hashInputs: [format, samplePoints, A, addressModeU, addressModeV, minFilter, offset], + }).map(({ coords, arrayIndex, depthRef, offset }) => { + return { + builtin: 'textureGatherCompare', + coordType: 'f', + coords, + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + depthRef, + offset, + }; + }); + const textureType = 'texture_depth_2d_array'; + 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('array_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegathercompare') .desc( ` -C: i32, u32 +A: i32, u32 -fn textureGatherCompare(t: texture_depth_cube_array, s: sampler_comparison, coords: vec3, array_index: C, depth_ref: f32) -> vec4 +fn textureGatherCompare(t: texture_depth_cube_array, s: sampler_comparison, coords: vec3, array_index: A, depth_ref: f32) -> vec4 Parameters: * t: The depth texture to read from @@ -73,17 +154,81 @@ Parameters: * depth_ref: The reference value to compare the sampled depth value against ` ) - .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]) - .combine('coords', generateCoordBoundaries(3)) - .combine('depth_ref', [-1 /* smaller ref */, 0 /* equal ref */, 1 /* larger ref */] 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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('A', ['i32', 'u32'] as const) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('compare', kCompareFunctions) ) - .unimplemented(); + .beforeAllSubcases(t => { + t.skipIfTextureViewDimensionNotSupported('cube-array'); + }) + .fn(async t => { + const { format, A, samplePoints, addressMode, minFilter, compare } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube-array'; + const size = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + addressModeW: addressMode, + compare, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; -g.test('sampled_array_2d_coords') + const calls: TextureCall[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + textureBuiltin: 'textureGatherCompare', + arrayIndex: { num: texture.depthOrArrayLayers / 6, type: A }, + depthRef: true, + hashInputs: [format, samplePoints, addressMode, minFilter], + }).map(({ coords, depthRef, arrayIndex }) => { + return { + builtin: 'textureGatherCompare', + arrayIndex, + arrayIndexType: A === 'i32' ? 'i' : 'u', + coordType: 'f', + coords, + depthRef, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = '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); + }); + +g.test('sampled_2d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegathercompare') .desc( ` @@ -103,16 +248,71 @@ 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('depth_ref', [-1 /* smaller ref */, 0 /* equal ref */, 1 /* larger ref */] as const) - .combine('offset', generateOffsets(2)) + .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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('C', ['i32', 'u32'] as const) + .combine('samplePoints', kSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('compare', kCompareFunctions) + .combine('offset', [false, true] as const) ) - .unimplemented(); + .fn(async t => { + const { format, C, samplePoints, addressMode, compare, minFilter, offset } = t.params; + + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 4, format }); + const descriptor: GPUTextureDescriptor = { + format, + size: { width, height }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + compare, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateTextureBuiltinInputs2D(50, { + method: samplePoints, + sampler, + descriptor, + offset, + depthRef: true, + hashInputs: [format, C, samplePoints, addressMode, minFilter, compare, offset], + }).map(({ coords, depthRef, offset }) => { + return { + builtin: 'textureGatherCompare', + coordType: 'f', + coords, + depthRef, + offset, + }; + }); + const textureType = 'texture_depth_2d'; + 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_3d_coords') +g.test('sampled_3d_coords') .specURL('https://www.w3.org/TR/WGSL/#texturegathercompare') .desc( ` @@ -125,10 +325,70 @@ Parameters: * depth_ref: The reference value to compare the sampled depth value against ` ) - .paramsSubcasesOnly(u => + .params(u => u - .combine('S', ['clamp-to-edge', 'repeat', 'mirror-repeat']) - .combine('coords', generateCoordBoundaries(3)) - .combine('depth_ref', [-1 /* smaller ref */, 0 /* equal ref */, 1 /* larger ref */] 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)) + .combine('minFilter', ['nearest', 'linear'] as const) + .beginSubcases() + .combine('samplePoints', kCubeSamplePointMethods) + .combine('addressMode', ['clamp-to-edge', 'repeat', 'mirror-repeat'] as const) + .combine('compare', kCompareFunctions) ) - .unimplemented(); + .fn(async t => { + const { format, samplePoints, addressMode, minFilter, compare } = t.params; + + const viewDimension: GPUTextureViewDimension = 'cube'; + const [width, height] = chooseTextureSize({ minSize: 8, minBlocks: 2, format, viewDimension }); + const depthOrArrayLayers = 6; + + const descriptor: GPUTextureDescriptor = { + format, + ...(t.isCompatibility && { textureBindingViewDimension: viewDimension }), + size: { width, height, depthOrArrayLayers }, + usage: GPUTextureUsage.COPY_DST | GPUTextureUsage.TEXTURE_BINDING, + }; + const { texels, texture } = await createTextureWithRandomDataAndGetTexels(t, descriptor); + const sampler: GPUSamplerDescriptor = { + addressModeU: addressMode, + addressModeV: addressMode, + addressModeW: addressMode, + compare, + minFilter, + magFilter: minFilter, + mipmapFilter: minFilter, + }; + + const calls: TextureCall[] = generateSamplePointsCube(50, { + method: samplePoints, + sampler, + descriptor, + depthRef: true, + textureBuiltin: 'textureGatherCompare', + hashInputs: [format, samplePoints, addressMode, minFilter, compare], + }).map(({ coords, depthRef }) => { + return { + builtin: 'textureGatherCompare', + coordType: 'f', + coords, + depthRef, + }; + }); + const viewDescriptor = { + dimension: viewDimension, + }; + const textureType = 'texture_depth_cube'; + 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 4da0c3dda5cf..c6de05725e12 100644 --- a/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts +++ b/src/webgpu/shader/execution/expression/call/builtin/texture_utils.ts @@ -535,6 +535,7 @@ const kTextureCallArgNames: readonly TextureCallArgKeys[] = [ 'mipLevel', 'ddx', 'ddy', + 'depthRef', 'offset', ] as const; @@ -544,6 +545,7 @@ export interface TextureCallArgs { mipLevel?: number; arrayIndex?: number; sampleIndex?: number; + depthRef?: number; ddx?: T; ddy?: T; offset?: T; @@ -551,6 +553,7 @@ export interface TextureCallArgs { export type TextureBuiltin = | 'textureGather' + | 'textureGatherCompare' | 'textureLoad' | 'textureSample' | 'textureSampleBaseClampToEdge' @@ -565,6 +568,10 @@ export interface TextureCall extends TextureCallArgs builtin === 'textureGatherCompare'; +const isBuiltinGather = (builtin: TextureBuiltin) => + builtin === 'textureGather' || builtin === 'textureGatherCompare'; + const s_u32 = new Uint32Array(1); const s_f32 = new Float32Array(s_u32.buffer); const s_i32 = new Int32Array(s_u32.buffer); @@ -605,6 +612,7 @@ function getCallArgType( case 'sampleIndex': assert(call.sampleIndexType !== undefined); return call.sampleIndexType; + case 'depthRef': case 'ddx': case 'ddy': return 'f'; @@ -718,6 +726,37 @@ function zeroValuePerTexelComponent(components: TexelComponent[]) { return out; } +const kSamplerFns: Record boolean> = { + never: (ref: number, v: number) => false, + less: (ref: number, v: number) => ref < v, + equal: (ref: number, v: number) => ref === v, + 'less-equal': (ref: number, v: number) => ref <= v, + greater: (ref: number, v: number) => ref > v, + 'not-equal': (ref: number, v: number) => ref !== v, + 'greater-equal': (ref: number, v: number) => ref >= v, + always: (ref: number, v: number) => true, +} as const; + +function applyCompare( + call: TextureCall, + sampler: GPUSamplerDescriptor | undefined, + components: TexelComponent[], + src: PerTexelComponent +): PerTexelComponent { + if (isBuiltinComparison(call.builtin)) { + assert(sampler !== undefined); + assert(call.depthRef !== undefined); + const out: PerTexelComponent = {}; + const compareFn = kSamplerFns[sampler.compare!]; + for (const component of components) { + out[component] = compareFn(call.depthRef, src[component]!) ? 1 : 0; + } + return out; + } else { + return src; + } +} + /** * Returns the expect value for a WGSL builtin texture function for a single * mip level @@ -769,6 +808,7 @@ export function softwareTextureReadMipLevel( switch (call.builtin) { case 'textureGather': + case 'textureGatherCompare': case 'textureSample': case 'textureSampleBaseClampToEdge': case 'textureSampleLevel': { @@ -799,7 +839,7 @@ export function softwareTextureReadMipLevel( const samples: { at: number[]; weight: number }[] = []; - const filter = call.builtin === 'textureGather' ? 'linear' : sampler?.minFilter ?? 'nearest'; + const filter = isBuiltinGather(call.builtin) ? 'linear' : sampler?.minFilter ?? 'nearest'; switch (filter) { case 'linear': { // 'p0' is the lower texel for 'at' @@ -909,7 +949,7 @@ export function softwareTextureReadMipLevel( unreachable(); } - if (call.builtin === 'textureGather') { + if (isBuiltinGather(call.builtin)) { const componentNdx = call.component ?? 0; assert(componentNdx >= 0 && componentNdx < 4); assert(samples.length === 4); @@ -920,23 +960,22 @@ export function softwareTextureReadMipLevel( ? wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize[0], sample.at as vec3) : applyAddressModesToCoords(addressMode, textureSize, sample.at); const v = load(c); - const rgba = convertPerTexelComponentToResultFormat(v, format); + const postV = applyCompare(call, sampler, rep.componentOrder, v); + const rgba = convertPerTexelComponentToResultFormat(postV, format); out[kRGBAComponents[i]] = rgba[component]; }); - return out; } const out: PerTexelComponent = {}; - const ss = []; for (const sample of samples) { const c = isCube ? wrapFaceCoordToCubeFaceAtEdgeBoundaries(textureSize[0], sample.at as vec3) : applyAddressModesToCoords(addressMode, textureSize, sample.at); const v = load(c); - ss.push(v); + const postV = applyCompare(call, sampler, rep.componentOrder, v); for (const component of rep.componentOrder) { - out[component] = (out[component] ?? 0) + v[component]! * sample.weight; + out[component] = (out[component] ?? 0) + postV[component]! * sample.weight; } } @@ -1225,7 +1264,7 @@ function getULPFromZeroForComponents( componentNdx?: number ): PerTexelComponent { const rep = kTexelRepresentationInfo[format]; - if (builtin === 'textureGather') { + if (isBuiltinGather(builtin)) { const out: PerTexelComponent = {}; const component = kRGBAComponents[componentNdx ?? 0]; const temp: PerTexelComponent = { R: 0, G: 0, B: 0, A: 1 }; @@ -1278,7 +1317,7 @@ export async function checkCallResults( // so if this is `textureGather` and component > 0 then there's nothing to check. if ( isDepthOrStencilTextureFormat(format) && - call.builtin === 'textureGather' && + isBuiltinGather(call.builtin) && call.component! > 0 ) { continue; @@ -1298,7 +1337,7 @@ export async function checkCallResults( // from the spec: https://gpuweb.github.io/gpuweb/#reading-depth-stencil // depth and stencil values are D, ?, ?, ? const rgbaComponentsToCheck = - call.builtin === 'textureGather' || !isDepthOrStencilTextureFormat(format) + isBuiltinGather(call.builtin) || !isDepthOrStencilTextureFormat(format) ? kRGBAComponents : kRComponent; @@ -2135,9 +2174,10 @@ async function identifySamplePoints( ) as EncodableTextureFormat; const rep = kTexelRepresentationInfo[format]; - const components = call.builtin === 'textureGather' ? kRGBAComponents : rep.componentOrder; - const convertResultAsAppropriate = - call.builtin === 'textureGather' ? (v: T) => v : convertResultFormatToTexelViewFormat; + const components = isBuiltinGather(call.builtin) ? kRGBAComponents : rep.componentOrder; + const convertResultAsAppropriate = isBuiltinGather(call.builtin) + ? (v: T) => v + : convertResultFormatToTexelViewFormat; // Identify all the texels that are sampled, and their weights. const sampledTexelWeights = new Map>(); @@ -2156,7 +2196,7 @@ async function identifySamplePoints( unclassifiedStack.push(setB); } - // See if any of the texels in setA were sampled. + // See if any of the texels in setA were sampled.0 const results = convertResultAsAppropriate( await run( range(mipLevelCount, mipLevel => @@ -2393,6 +2433,7 @@ type TextureBuiltinInputArgs = { sampleIndex?: RangeDef; arrayIndex?: RangeDef; component?: boolean; + depthRef?: boolean; offset?: boolean; hashInputs: (number | string | boolean)[]; }; @@ -2419,6 +2460,7 @@ function generateTextureBuiltinInputsImpl( arrayIndex?: number; offset?: T; component?: number; + depthRef?: number; }[] { const { method, descriptor } = args; const dimension = descriptor.dimension ?? '2d'; @@ -2473,9 +2515,7 @@ function generateTextureBuiltinInputsImpl( // MacOS, M1 Mac: 256 const kSubdivisionsPerTexel = 4; const avoidEdgeCase = - !args.sampler || - args.sampler.minFilter === 'nearest' || - args.textureBuiltin === 'textureGather'; + !args.sampler || args.sampler.minFilter === 'nearest' || isBuiltinGather(args.textureBuiltin!); const edgeRemainder = args.textureBuiltin === 'textureGather' ? kSubdivisionsPerTexel / 2 : 0; const numComponents = isDepthOrStencilTextureFormat(descriptor.format) ? 1 : 4; return coords.map((c, i) => { @@ -2502,6 +2542,7 @@ function generateTextureBuiltinInputsImpl( mipLevel, sampleIndex: args.sampleIndex ? makeRangeValue(args.sampleIndex, i, 1) : undefined, arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined, + depthRef: args.depthRef ? makeRangeValue({ num: 1, type: 'f32' }, i, 5) : undefined, offset: args.offset ? (coords.map((_, j) => makeIntHashValueRepeatable(-8, 8, i, 3 + j)) as T) : undefined, @@ -2708,6 +2749,7 @@ export function generateSamplePointsCube( arrayIndex?: number; offset?: undefined; component?: number; + depthRef?: number; }[] { const { method, descriptor } = args; const mipLevelCount = descriptor.mipLevelCount ?? 1; @@ -2906,10 +2948,8 @@ export function generateSamplePointsCube( // const kSubdivisionsPerTexel = 4; const avoidEdgeCase = - !args.sampler || - args.sampler.minFilter === 'nearest' || - args.textureBuiltin === 'textureGather'; - const edgeRemainder = args.textureBuiltin === 'textureGather' ? kSubdivisionsPerTexel / 2 : 0; + !args.sampler || args.sampler.minFilter === 'nearest' || isBuiltinGather(args.textureBuiltin!); + const edgeRemainder = isBuiltinGather(args.textureBuiltin!) ? kSubdivisionsPerTexel / 2 : 0; return coords.map((c, i) => { const mipLevel = args.mipLevel @@ -2949,6 +2989,7 @@ export function generateSamplePointsCube( coords, mipLevel, arrayIndex: args.arrayIndex ? makeRangeValue(args.arrayIndex, i, 2) : undefined, + depthRef: args.depthRef ? makeRangeValue({ num: 1, type: 'f32' }, i, 5) : undefined, component: args.component ? makeIntHashValue(0, 4, i, 4) : undefined, }; }); @@ -3054,6 +3095,8 @@ function buildBinnedCalls(calls: TextureCall[]) { ? prototype.arrayIndexType! : name === 'sampleIndex' ? prototype.sampleIndexType! + : name === 'depthRef' + ? 'f' : prototype.coordType; args.push(`args.${name}`); fields.push(`@align(16) ${name} : ${wgslTypeFor(value, type)}`); @@ -3127,6 +3170,8 @@ export function describeTextureCall(call: TextureCall< args.push(`${name}: ${wgslExprFor(value, call.arrayIndexType!)}`); } else if (name === 'sampleIndex') { args.push(`${name}: ${wgslExprFor(value, call.sampleIndexType!)}`); + } else if (name === 'depthRef') { + args.push(`${name}: ${wgslExprFor(value, 'f')}`); } else { args.push(`${name}: ${wgslExpr(value)}`); } @@ -3204,16 +3249,20 @@ export async function doTextureCalls( }); t.device.queue.writeBuffer(dataBuffer, 0, new Uint32Array(data)); - const { resultType, resultFormat, componentType } = - calls[0].builtin === 'textureGather' - ? getTextureFormatTypeInfo(format) - : gpuTexture instanceof GPUExternalTexture - ? ({ resultType: 'vec4f', resultFormat: 'rgba32float', componentType: 'f32' } as const) - : textureType.includes('depth') - ? ({ resultType: 'f32', resultFormat: 'rgba32float', componentType: 'f32' } as const) - : getTextureFormatTypeInfo(format); + const builtin = calls[0].builtin; + const isCompare = isBuiltinComparison(builtin); + + const { resultType, resultFormat, componentType } = isBuiltinGather(builtin) + ? getTextureFormatTypeInfo(format) + : gpuTexture instanceof GPUExternalTexture + ? ({ resultType: 'vec4f', resultFormat: 'rgba32float', componentType: 'f32' } as const) + : textureType.includes('depth') + ? ({ resultType: 'f32', resultFormat: 'rgba32float', componentType: 'f32' } as const) + : getTextureFormatTypeInfo(format); const returnType = `vec4<${componentType}>`; + const samplerType = isCompare ? 'sampler_comparison' : 'sampler'; + const rtWidth = 256; const renderTarget = t.createTextureTracked({ format: resultFormat, @@ -3238,7 +3287,7 @@ fn vs_main(@builtin(vertex_index) vertex_index : u32) -> @builtin(position) vec4 } @group(0) @binding(0) var T : ${textureType}; -${sampler ? '@group(0) @binding(1) var S : sampler' : ''}; +${sampler ? `@group(0) @binding(1) var S : ${samplerType}` : ''}; @group(0) @binding(2) var data : Data; @fragment @@ -3329,7 +3378,7 @@ ${body} binding: 1, visibility: GPUShaderStage.FRAGMENT, sampler: { - type: isFiltering ? 'filtering' : 'non-filtering', + type: isCompare ? 'comparison' : isFiltering ? 'filtering' : 'non-filtering', }, }); } From 072d7b682e12996f3ac733ce55530aa619c54cc4 Mon Sep 17 00:00:00 2001 From: alan-baker Date: Fri, 23 Aug 2024 14:26:14 -0400 Subject: [PATCH 18/99] Tests for subgroup builtin values in fragment shaders (#3915) * Tests subgroup_size and subgroup_invocation_id in fragment shaders --- .../shader_io/fragment_builtins.spec.ts | 394 +++++++++++++++++- 1 file changed, 388 insertions(+), 6 deletions(-) diff --git a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts index ffd58976fc88..6701315d05c1 100644 --- a/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts +++ b/src/webgpu/shader/execution/shader_io/fragment_builtins.spec.ts @@ -20,14 +20,17 @@ is evaluated per-fragment or per-sample. With @interpolate(, sample) or usage of import { makeTestGroup } from '../../../../common/framework/test_group.js'; import { ErrorWithExtra, assert, range, unreachable } from '../../../../common/util/util.js'; import { InterpolationSampling, InterpolationType } from '../../../constants.js'; -import { GPUTest } from '../../../gpu_test.js'; +import { kTextureFormatInfo } from '../../../format_info.js'; +import { GPUTest, TextureTestMixin } from '../../../gpu_test.js'; import { getProvokingVertexForFlatInterpolationEitherSampling } from '../../../inter_stage.js'; import { getMultisampleFragmentOffsets } from '../../../multisample_info.js'; -import { dotProduct, subtractVectors } from '../../../util/math.js'; +import { dotProduct, subtractVectors, align } from '../../../util/math.js'; import { TexelView } from '../../../util/texture/texel_view.js'; import { findFailedPixels } from '../../../util/texture/texture_ok.js'; -export const g = makeTestGroup(GPUTest); +class FragmentBuiltinTest extends TextureTestMixin(GPUTest) {} + +export const g = makeTestGroup(FragmentBuiltinTest); const s_deviceToPipelineMap = new WeakMap< GPUDevice, @@ -589,7 +592,7 @@ async function renderFragmentShaderInputsTo4TexturesAndReadbackValues( struct FragmentIn { @builtin(position) position: vec4f, - @location(0) @interpolate(${interpolate}) interpolatedValue: vec4f, +@location(0) @interpolate(${interpolate}) interpolatedValue: vec4f, ${fragInCode} }; @@ -1424,6 +1427,385 @@ g.test('inputs,sample_mask') ); }); -g.test('subgroup_size').unimplemented(); +const kSizes = [ + [15, 15], + [16, 16], + [17, 17], + [19, 13], + [13, 10], + [111, 2], + [2, 111], + [35, 2], + [2, 35], + [53, 13], + [13, 53], +] as const; + +/** + * @returns The population count of input. + * + * @param input Treated as an unsigned 32-bit integer + */ +function popcount(input: number): number { + let n = input; + n = n - ((n >> 1) & 0x55555555); + n = (n & 0x33333333) + ((n >> 2) & 0x33333333); + return (((n + (n >> 4)) & 0xf0f0f0f) * 0x1010101) >> 24; +} + +/** + * Checks subgroup_size builtin value consistency. + * + * The builtin subgroup_size is not assumed to be uniform in fragment shaders. + * Therefore, this function checks the value is a power of two within the device + * limits and that the ballot size is less than the stated size. + * @param data An array of vec4u that contains (per texel): + * * builtin value + * * ballot size + * * comparison to other invocations + * * 0 + * @param format The texture format for data + * @param min The minimum subgroup size from the device + * @param max The maximum subgroup size from the device + * @param width The width of the framebuffer + * @param height The height of the framebuffer + */ +function checkSubgroupSizeConsistency( + data: Uint32Array, + format: GPUTextureFormat, + min: number, + max: number, + width: number, + height: number +): Error | undefined { + const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format]; + const blocksPerRow = width / blockWidth; + // Image copies require bytesPerRow to be a multiple of 256. + const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256); + const uintsPerRow = bytesPerRow / 4; + const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4; + + for (let row = 0; row < height; row++) { + for (let col = 0; col < width; col++) { + const offset = uintsPerRow * row + col * uintsPerTexel; + const builtinSize = data[offset]; + const ballotSize = data[offset + 1]; + const comparison = data[offset + 2]; + if (builtinSize === 0) { + continue; + } + + if (popcount(builtinSize) !== 1) { + return new Error(`Subgroup size '${builtinSize}' is not a power of two`); + } + + if (builtinSize < min) { + return new Error(`Subgroup size '${builtinSize}' is less than minimum '${min}'`); + } + if (max < builtinSize) { + return new Error(`Subgroup size '${builtinSize}' is greater than maximum '${max}'`); + } + + if (builtinSize < ballotSize) { + return new Error(`Inconsistent subgroup ballot size +- icoord: (${row}, ${col}) +- expected: ${builtinSize} +- got: ${ballotSize}`); + } + + if (comparison !== 1) { + return new Error(`Not all invocations in subgroup have same view of the size +- icoord: (${row}, ${col})`); + } + } + } + + return undefined; +} + +/** + * Runs a subgroup builtin test for fragment shaders + * + * This test draws a full screen in 2 separate draw calls (half screen each). + * Results are checked for each draw. + * @param t The base test + * @param format The framebuffer format + * @param fsShader The fragment shader with the following interface: + * Location 0 output is framebuffer with format + * Group 0 binding 0 is a u32 sized data + * @param width The framebuffer width + * @param height The framebuffer height + * @param checker A functor to check the framebuffer values + */ +async function runSubgroupTest( + t: FragmentBuiltinTest, + format: GPUTextureFormat, + fsShader: string, + width: number, + height: number, + checker: (data: Uint32Array) => Error | undefined +) { + const vsShader = ` +@vertex +fn vsMain(@builtin(vertex_index) index : u32) -> @builtin(position) vec4f { + const vertices = array( + vec2(-1, -1), vec2(-1, 1), vec2( 1, 1), + vec2(-1, -1), vec2( 1, -1), vec2( 1, 1), + ); + return vec4f(vec2f(vertices[index]), 0, 1); +}`; + + const pipeline = t.device.createRenderPipeline({ + layout: 'auto', + vertex: { + module: t.device.createShaderModule({ code: vsShader }), + }, + fragment: { + module: t.device.createShaderModule({ code: fsShader }), + targets: [{ format }], + }, + primitive: { + topology: 'triangle-list', + }, + }); + + const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format]; + assert(bytesPerBlock !== undefined); + + const blocksPerRow = width / blockWidth; + const blocksPerColumn = height / blockHeight; + const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256); + const byteLength = bytesPerRow * blocksPerColumn; + const uintLength = byteLength / 4; + + const buffer = t.makeBufferWithContents( + new Uint32Array([1]), + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const bg = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer, + }, + }, + ], + }); -g.test('subgroup_invocation_id').unimplemented(); + for (let i = 0; i < 2; i++) { + const framebuffer = t.createTextureTracked({ + size: [width, height], + usage: + GPUTextureUsage.COPY_SRC | + GPUTextureUsage.COPY_DST | + GPUTextureUsage.RENDER_ATTACHMENT | + GPUTextureUsage.TEXTURE_BINDING, + format, + }); + + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: framebuffer.createView(), + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bg); + pass.draw(3, 1, i); + pass.end(); + t.queue.submit([encoder.finish()]); + + const buffer = t.copyWholeTextureToNewBufferSimple(framebuffer, 0); + const readback = await t.readGPUBufferRangeTyped(buffer, { + srcByteOffset: 0, + type: Uint32Array, + typedLength: uintLength, + method: 'copy', + }); + const data: Uint32Array = readback.data; + + t.expectOK(checker(data)); + } +} + +g.test('subgroup_size') + .desc('Tests subgroup_size values') + .params(u => + u + .combine('size', kSizes) + .beginSubcases() + .combineWithParams([{ format: 'rgba32uint' }] as const) + ) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(async t => { + interface SubgroupLimits extends GPUSupportedLimits { + minSubgroupSize: number; + maxSubgroupSize: number; + } + const { minSubgroupSize, maxSubgroupSize } = t.device.limits as SubgroupLimits; + + const fsShader = ` +enable subgroups; + +const width = ${t.params.size[0]}; +const height = ${t.params.size[1]}; + +@group(0) @binding(0) var var for_layout : u32; + +@fragment +fn fsMain( + _ = for_layout; + + @builtin(position) pos : vec4f, + @builtin(subgroup_size) sg_size : u32, +) -> @location(0) vec4u { + let ballot = countOneBits(subgroupBallot(true)); + let ballotSize = ballot.x + ballot.y + ballot.z + ballot.w; + + // Do all invocations in the subgroup see the same subgroup size? + let firstSize = subgroupBroadcast(sg_size, 0); + let compareBallot = countOneBits(subgroupBallot(firstSize == sg_size)); + let compareSize = compareBallot.x + compareBallot.y + compareBallot.z + compareBallot.w; + let sameSize = select(0u, 1u, compareSize == ballotSize); + + return vec4u(sg_size, ballotSize, sameSize, 0); +}`; + + await runSubgroupTest( + t, + t.params.format, + fsShader, + t.params.size[0], + t.params.size[1], + (data: Uint32Array) => { + return checkSubgroupSizeConsistency( + data, + t.params.format, + minSubgroupSize, + maxSubgroupSize, + t.params.size[0], + t.params.size[1] + ); + } + ); + }); + +/** + * Checks subgroup_invocation_id value consistency + * + * Very little uniformity is expected for subgroup_invocation_id. + * This function checks that all ids are less than the subgroup size + * and no id is repeated. + * @param data An array of vec4u that contains (per texel): + * * subgroup_invocation_id + * * ballot size + * * non-zero ID unique to each subgroup + * * 0 + * @param format The texture format of data + * @param width The width of the framebuffer + * @param height The height of the framebuffer + */ +function checkSubgroupInvocationIdConsistency( + data: Uint32Array, + format: GPUTextureFormat, + width: number, + height: number +): Error | undefined { + const { blockWidth, blockHeight, bytesPerBlock } = kTextureFormatInfo[format]; + const blocksPerRow = width / blockWidth; + const bytesPerRow = align(blocksPerRow * (bytesPerBlock ?? 1), 256); + const uintsPerRow = bytesPerRow / 4; + const uintsPerTexel = (bytesPerBlock ?? 1) / blockWidth / blockHeight / 4; + + const mappings = new Map(); + for (let row = 0; row < height; row++) { + for (let col = 0; col < width; col++) { + const offset = uintsPerRow * row + col * uintsPerTexel; + const id = data[offset]; + const size = data[offset + 1]; + const repId = data[offset + 2]; + + if (repId === 0) { + continue; + } + + if (size < id) { + return new Error( + `Invocation id '${id}' is greater than subgroup size '${size}' for (${row}, ${col})` + ); + } + + let v = mappings.get(repId) ?? 0n; + const mask = 1n << BigInt(id); + if ((mask & v) !== 0n) { + return new Error(`Multiple invocations with id '${id}' in subgroup '${repId}'`); + } + v |= mask; + mappings.set(repId, v); + } + } + + return undefined; +} + +g.test('subgroup_invocation_id') + .desc('Tests subgroup_invocation_id built-in value') + .params(u => + u + .combine('size', kSizes) + .beginSubcases() + .combineWithParams([{ format: 'rgba32uint' }] as const) + ) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(async t => { + const fsShader = ` +enable subgroups; + +const width = ${t.params.size[0]}; +const height = ${t.params.size[1]}; + +@group(0) @binding(0) var counter : atomic; + +@fragment +fn fsMain( + @builtin(position) pos : vec4f, + @builtin(subgroup_invocation_id) id : u32, + @builtin(subgroup_size) sg_size : u32, +) -> @location(0) vec4u { + let ballot = countOneBits(subgroupBallot(true)); + let ballotSize = ballot.x + ballot.y + ballot.z + ballot.w; + + // Generate representative id for this subgroup. + var repId = atomicAdd(&counter, 1); + repId = subgroupBroadcast(repId, 0); + + return vec4u(id, ballotSize, repId, 0); +}`; + + await runSubgroupTest( + t, + t.params.format, + fsShader, + t.params.size[0], + t.params.size[1], + (data: Uint32Array) => { + return checkSubgroupInvocationIdConsistency( + data, + t.params.format, + t.params.size[0], + t.params.size[1] + ); + } + ); + }); From 9bde59301a5ec19bfd77bb72e5e9230f74ebaae1 Mon Sep 17 00:00:00 2001 From: David Neto Date: Fri, 23 Aug 2024 16:45:08 -0400 Subject: [PATCH 19/99] Add Chrome OT token for "WebGPU Subgroups Feature" (#3895) Remove the OT tokens for Chrome's initial WebGPU feature launch. crbug.com/358117283 --- src/common/tools/dev_server.ts | 4 ++-- standalone/index.html | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/src/common/tools/dev_server.ts b/src/common/tools/dev_server.ts index 1d1313e4f51c..a14ccb54961f 100644 --- a/src/common/tools/dev_server.ts +++ b/src/common/tools/dev_server.ts @@ -107,8 +107,8 @@ const app = express(); // Send Chrome Origin Trial tokens app.use((_req, res, next) => { res.header('Origin-Trial', [ - // Token for http://localhost:8080 - 'AvyDIV+RJoYs8fn3W6kIrBhWw0te0klraoz04mw/nPb8VTus3w5HCdy+vXqsSzomIH745CT6B5j1naHgWqt/tw8AAABJeyJvcmlnaW4iOiJodHRwOi8vbG9jYWxob3N0OjgwODAiLCJmZWF0dXJlIjoiV2ViR1BVIiwiZXhwaXJ5IjoxNjYzNzE4Mzk5fQ==', + // "WebGPU Subgroups Feature" token for http://localhost:8080 + 'AkMLfHisU+Fsbpi9g6tfKSZF4ngpsmjW4Oai360fUvZE2rgSPZDWSWb8ryrliJX5HR/Rw0yig0ir9el2hrnODwcAAABaeyJvcmlnaW4iOiJodHRwOi8vbG9jYWxob3N0OjgwODAiLCJmZWF0dXJlIjoiV2ViR1BVU3ViZ3JvdXBzRmVhdHVyZXMiLCJleHBpcnkiOjE3Mzk5MjMxOTl9', ]); next(); }); diff --git a/standalone/index.html b/standalone/index.html index d087d6584cd9..eb58b6102404 100644 --- a/standalone/index.html +++ b/standalone/index.html @@ -13,8 +13,8 @@ - - + +