From fe65191fe06fcf866082556b6fd8d0d2b0e18d30 Mon Sep 17 00:00:00 2001 From: alan-baker Date: Fri, 29 Nov 2024 14:14:13 -0500 Subject: [PATCH] Uniformity tests for subgroup built-in functions (#4064) * Built-in values already tested * Tests call site requirements and parameter requirements --- src/webgpu/listing_meta.json | 2 + .../validation/uniformity/uniformity.spec.ts | 145 ++++++++++++++++++ 2 files changed, 147 insertions(+) diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 81f2449c9e7c..fb5e451b32a0 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -2918,6 +2918,7 @@ "webgpu:shader,validation,types,textures:texel_formats,as_value:*": { "subcaseMS": 0.518 }, "webgpu:shader,validation,types,textures:texel_formats:*": { "subcaseMS": 1707.432 }, "webgpu:shader,validation,types,vector:vector:*": { "subcaseMS": 1.295 }, + "webgpu:shader,validation,uniformity,uniformity:basics,subgroups:*": { "subcaseMS": 5413.204 }, "webgpu:shader,validation,uniformity,uniformity:basics:*": { "subcaseMS": 1.467 }, "webgpu:shader,validation,uniformity,uniformity:binary_expressions:*": { "subcaseMS": 1.758 }, "webgpu:shader,validation,uniformity,uniformity:compute_builtin_values:*": { "subcaseMS": 2.500 }, @@ -2927,6 +2928,7 @@ "webgpu:shader,validation,uniformity,uniformity:functions:*": { "subcaseMS": 1.303 }, "webgpu:shader,validation,uniformity,uniformity:pointers:*": { "subcaseMS": 1.738 }, "webgpu:shader,validation,uniformity,uniformity:short_circuit_expressions:*": { "subcaseMS": 1.401 }, + "webgpu:shader,validation,uniformity,uniformity:subgroups,parameters:*": { "subcaseMS": 81.282 }, "webgpu:shader,validation,uniformity,uniformity:unary_expressions:*": { "subcaseMS": 1.279 }, "webgpu:util,texture,color_space_conversions:util_matches_2d_canvas:*": { "subcaseMS": 1.001 }, "webgpu:util,texture,texel_data:float_texel_data_in_shader:*": { "subcaseMS": 2.042 }, diff --git a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts index da34d593b620..b3b3cdb1db43 100644 --- a/src/webgpu/shader/validation/uniformity/uniformity.spec.ts +++ b/src/webgpu/shader/validation/uniformity/uniformity.spec.ts @@ -143,6 +143,35 @@ function generateOp(op: string): string { case 'fwidthFine': { return `let x = ${op}(0);\n`; } + case 'subgroupAdd': + case 'subgroupInclusiveAdd': + case 'subgroupExclusiveAdd': + case 'subgroupMul': + case 'subgroupInclusiveMul': + case 'subgroupExclusiveMul': + case 'subgroupMax': + case 'subgroupMin': + case 'subgroupAnd': + case 'subgroupOr': + case 'subgroupXor': + case 'subgroupBroadcastFirst': + case 'quadSwapX': + case 'quadSwapY': + case 'quadSwapDiagonal': + return `let x = ${op}(0);\n`; + case 'subgroupAll': + case 'subgroupAny': + case 'subgroupBallot': + return `let x = ${op}(false);\n`; + case 'subgroupElect': + return `let x = ${op}();\n`; + case 'subgroupBroadcast': + case 'subgroupShuffle': + case 'subgroupShuffleUp': + case 'subgroupShuffleDown': + case 'subgroupShuffleXor': + case 'quadBroadcast': + return `let x = ${op}(0, 0);\n`; default: { unreachable(`Unhandled op`); } @@ -246,6 +275,95 @@ g.test('basics') t.expectCompileResult(t.params.expectation, code); }); +const kSubgroupOps = [ + 'subgroupAdd', + 'subgroupInclusiveAdd', + 'subgroupExclusiveAdd', + 'subgroupMul', + 'subgroupInclusiveMul', + 'subgroupExclusiveMul', + 'subgroupMax', + 'subgroupMin', + 'subgroupAll', + 'subgroupAny', + 'subgroupAnd', + 'subgroupOr', + 'subgroupXor', + 'subgroupBallot', + 'subgroupElect', + 'subgroupBroadcast', + 'subgroupBroadcastFirst', + 'subgroupShuffle', + 'subgroupShuffleUp', + 'subgroupShuffleDown', + 'subgroupShuffleXor', + 'quadBroadcast', + 'quadSwapX', + 'quadSwapY', + 'quadSwapDiagonal', +] as const; + +g.test('basics,subgroups') + .desc(`Test subgroup operations in simple uniform or non-uniform control flow.`) + .params(u => + u + .combine('statement', ['if', 'for', 'while', 'switch'] as const) + .beginSubcases() + .combineWithParams(kConditions) + .combine('op', kSubgroupOps) + .combine('stage', ['compute', 'fragment'] as const) + ) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + let code = ` + enable subgroups; + + @group(0) @binding(0) var s : sampler; + @group(0) @binding(1) var s_comp : sampler_comparison; + @group(0) @binding(2) var tex : texture_2d; + @group(0) @binding(3) var tex_depth : texture_depth_2d; + + @group(1) @binding(0) var ro_buffer : array; + @group(1) @binding(1) var rw_buffer : array; + @group(1) @binding(2) var uniform_buffer : vec4; + + @group(2) @binding(0) var ro_storage_texture : texture_storage_2d; + @group(2) @binding(1) var rw_storage_texture : texture_storage_2d; + + var priv_var : array = array(0,0,0,0); + + const c = false; + override o : f32; +`; + + if (t.params.stage === 'compute') { + code += `var wg : f32;\n`; + code += ` @workgroup_size(16, 1, 1)`; + } + code += `@${t.params.stage}`; + code += `\nfn main(`; + if (t.params.stage === 'compute') { + code += `@builtin(global_invocation_id) p : vec3`; + } else { + code += `@builtin(position) p : vec4`; + } + code += `) { + let u_let = uniform_buffer.x; + let n_let = rw_buffer[0]; + var u_f = uniform_buffer.z; + var n_f = rw_buffer[1]; + `; + + // Simple control statement containing the op. + code += generateConditionalStatement(t.params.statement, t.params.cond, t.params.op); + + code += `\n}\n`; + + t.expectCompileResult(t.params.expectation, code); + }); + const kFragmentBuiltinValues = [ { builtin: `position`, @@ -2693,3 +2811,30 @@ fn main() { } t.expectCompileResult(res, code); }); + +g.test('subgroups,parameters') + .desc('Test subgroup operations that require a uniform parameter') + .params(u => + u + .combine('op', ['subgroupShuffleUp', 'subgroupShuffleDown', 'subgroupShuffleXor'] as const) + .combine('uniform', [false, true] as const) + ) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; + +var non_uniform : u32 = 0; + +@group(0) @binding(0) +var uniform : u32; + +@compute @workgroup_size(16,1,1) +fn main() { + let x = ${t.params.op}(non_uniform, ${t.params.uniform ? 'uniform' : 'non_uniform'}); +}`; + + t.expectCompileResult(t.params.uniform, wgsl); + });