From 13d5f1781df552ab4b2f37a36a360c709f7d7343 Mon Sep 17 00:00:00 2001 From: alan-baker Date: Tue, 27 Aug 2024 23:00:04 -0400 Subject: [PATCH] Remaining subgroup validation tests (#3920) * Remaining subgroup validation tests * Adds must_use tests to ballot * Adds const id requirement to broadcast * Adds tests for other builtins: * broadcast first, * elect, * min, max, * and, or, xor, * shuffle, shuffle xor, shuffle up, shuffle down * quad broadcast * quad swap x, quad swap y, quad swap diagonal * Add enable validation tests to broadcast and ballot --- src/webgpu/listing_meta.json | 61 ++++ .../call/builtin/quadBroadcast.spec.ts | 286 ++++++++++++++++++ .../expression/call/builtin/quadSwap.spec.ts | 227 ++++++++++++++ .../call/builtin/subgroupAnyAll.spec.ts | 186 ++++++++++++ .../call/builtin/subgroupBallot.spec.ts | 35 ++- .../call/builtin/subgroupBitwise.spec.ts | 204 +++++++++++++ .../call/builtin/subgroupBroadcast.spec.ts | 90 ++++++ .../builtin/subgroupBroadcastFirst.spec.ts | 210 +++++++++++++ .../call/builtin/subgroupElect.spec.ts | 175 +++++++++++ .../call/builtin/subgroupMinMax.spec.ts | 227 ++++++++++++++ .../call/builtin/subgroupShuffle.spec.ts | 262 ++++++++++++++++ 11 files changed, 1962 insertions(+), 1 deletion(-) create mode 100644 src/webgpu/shader/validation/expression/call/builtin/quadBroadcast.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/quadSwap.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupAnyAll.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupBitwise.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcastFirst.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupElect.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupMinMax.spec.ts create mode 100644 src/webgpu/shader/validation/expression/call/builtin/subgroupShuffle.spec.ts diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index 0662ee913898..0290bae1808a 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -2234,6 +2234,22 @@ "webgpu:shader,validation,expression,call,builtin,pow:invalid_argument:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,expression,call,builtin,pow:must_use:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,expression,call,builtin,pow:values:*": { "subcaseMS": 1.000 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:data_type:*": { "subcaseMS": 39.783 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:early_eval:*": { "subcaseMS": 63.825 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:id_constness:*": { "subcaseMS": 15.347 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:id_type:*": { "subcaseMS": 26.268 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:must_use:*": { "subcaseMS": 41.658 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:requires_subgroups:*": { "subcaseMS": 42.565 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:requires_subgroups_f16:*": { "subcaseMS": 44.998 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:return_type:*": { "subcaseMS": 363.607 }, + "webgpu:shader,validation,expression,call,builtin,quadBroadcast:stage:*": { "subcaseMS": 3.050 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:data_type:*": { "subcaseMS": 89.379 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:early_eval:*": { "subcaseMS": 108.243 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:must_use:*": { "subcaseMS": 5.557 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:requires_subgroups:*": { "subcaseMS": 113.624 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:requires_subgroups_f16:*": { "subcaseMS": 12.712 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:return_type:*": { "subcaseMS": 1424.551 }, + "webgpu:shader,validation,expression,call,builtin,quadSwap:stage:*": { "subcaseMS": 7.664 }, "webgpu:shader,validation,expression,call,builtin,quantizeToF16:args:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,expression,call,builtin,quantizeToF16:must_use:*": { "subcaseMS": 1.000 }, "webgpu:shader,validation,expression,call,builtin,quantizeToF16:values:*": { "subcaseMS": 1.000 }, @@ -2292,22 +2308,67 @@ "webgpu:shader,validation,expression,call,builtin,subgroupAdd:must_use:*": { "subcaseMS": 62.933 }, "webgpu:shader,validation,expression,call,builtin,subgroupAdd:return_type:*": { "subcaseMS": 363.546 }, "webgpu:shader,validation,expression,call,builtin,subgroupAdd:stage:*": { "subcaseMS": 3.536 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:data_type:*": { "subcaseMS": 57.943 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:early_eval:*": { "subcaseMS": 173.714 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:must_use:*": { "subcaseMS": 4.592 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:requires_subgroups:*": { "subcaseMS": 73.866 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:return_type:*": { "subcaseMS": 39.388 }, + "webgpu:shader,validation,expression,call,builtin,subgroupAnyAll:stage:*": { "subcaseMS": 6.862 }, "webgpu:shader,validation,expression,call,builtin,subgroupBallot:data_type:*": { "subcaseMS": 115.557 }, "webgpu:shader,validation,expression,call,builtin,subgroupBallot:early_eval:*": { "subcaseMS": 52.992 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBallot:must_use:*": { "subcaseMS": 39.441 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBallot:requires_subgroups:*": { "subcaseMS": 36.819 }, "webgpu:shader,validation,expression,call,builtin,subgroupBallot:return_type:*": { "subcaseMS": 22.381 }, "webgpu:shader,validation,expression,call,builtin,subgroupBallot:stage:*": { "subcaseMS": 3.712 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:data_type:*": { "subcaseMS": 94.072 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:early_eval:*": { "subcaseMS": 569.598 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:must_use:*": { "subcaseMS": 6.172 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:requires_subgroups:*": { "subcaseMS": 108.478 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:return_type:*": { "subcaseMS": 1430.736 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBitwise:stage:*": { "subcaseMS": 11.858 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:data_type:*": { "subcaseMS": 97.991 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:early_eval:*": { "subcaseMS": 1.254 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:id_constness:*": { "subcaseMS": 7.026 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:id_type:*": { "subcaseMS": 24.703 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:must_use:*": { "subcaseMS": 232.030 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:requires_subgroups:*": { "subcaseMS": 47.231 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:requires_subgroups_f16:*": { "subcaseMS": 38.503 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:return_type:*": { "subcaseMS": 496.031 }, "webgpu:shader,validation,expression,call,builtin,subgroupBroadcast:stage:*": { "subcaseMS": 3.715 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:data_type:*": { "subcaseMS": 32.168 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:early_eval:*": { "subcaseMS": 57.922 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:must_use:*": { "subcaseMS": 36.296 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:requires_subgroups:*": { "subcaseMS": 42.522 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:requires_subgroups_f16:*": { "subcaseMS": 47.111 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:return_type:*": { "subcaseMS": 402.558 }, + "webgpu:shader,validation,expression,call,builtin,subgroupBroadcastFirst:stage:*": { "subcaseMS": 2.869 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:data_type:*": { "subcaseMS": 72.441 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:early_eval:*": { "subcaseMS": 56.115 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:must_use:*": { "subcaseMS": 32.820 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:requires_subgroups:*": { "subcaseMS": 35.595 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:return_type:*": { "subcaseMS": 22.712 }, + "webgpu:shader,validation,expression,call,builtin,subgroupElect:stage:*": { "subcaseMS": 3.790 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:data_type:*": { "subcaseMS": 64.143 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:early_eval:*": { "subcaseMS": 551.671 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:must_use:*": { "subcaseMS": 4.403 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:requires_subgroups:*": { "subcaseMS": 87.208 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:requires_subgroups_f16:*": { "subcaseMS": 25.190 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:return_type:*": { "subcaseMS": 911.454 }, + "webgpu:shader,validation,expression,call,builtin,subgroupMinMax:stage:*": { "subcaseMS": 6.395 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:data_type:*": { "subcaseMS": 45.396 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:early_eval:*": { "subcaseMS": 56.571 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:invalid_types:*": { "subcaseMS": 91.040 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:must_use:*": { "subcaseMS": 39.041 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:return_type:*": { "subcaseMS": 549.172 }, "webgpu:shader,validation,expression,call,builtin,subgroupMul:stage:*": { "subcaseMS": 4.489 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:data_type:*": { "subcaseMS": 115.093 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:early_eval:*": { "subcaseMS": 110.489 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:must_use:*": { "subcaseMS": 7.628 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:param2_type:*": { "subcaseMS": 88.305 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:requires_subgroups:*": { "subcaseMS": 102.779 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:requires_subgroups_f16:*": { "subcaseMS": 13.121 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:return_type:*": { "subcaseMS": 1930.309 }, + "webgpu:shader,validation,expression,call,builtin,subgroupShuffle:stage:*": { "subcaseMS": 9.527 }, "webgpu:shader,validation,expression,call,builtin,tan:args:*": { "subcaseMS": 43.560 }, "webgpu:shader,validation,expression,call,builtin,tan:must_use:*": { "subcaseMS": 5.401 }, "webgpu:shader,validation,expression,call,builtin,tan:values:*": { "subcaseMS": 0.350 }, diff --git a/src/webgpu/shader/validation/expression/call/builtin/quadBroadcast.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/quadBroadcast.spec.ts new file mode 100644 index 000000000000..6988f17b9ede --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/quadBroadcast.spec.ts @@ -0,0 +1,286 @@ +export const description = ` +Validation tests for quadBroadcast +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { + isConvertible, + Type, + elementTypeOf, + kAllScalarsAndVectors, +} from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = quadBroadcast(0, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = quadBroadcast(0h, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kArgumentTypes = objectsToRecord(kAllScalarsAndVectors); + +const kStages: Record = { + constant: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = quadBroadcast(0, 0); +}`, + override: ` +enable subgroups; +override o = quadBroadcast(0, 0);`, + runtime: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = quadBroadcast(0, 0); +}`, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage]; + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}quadBroadcast(0, 0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kArgumentTypes))) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kArgumentTypes[t.params.type]; + if (type.requiresF16()) { + features.push('subgroups-f16' as GPUFeatureName); + features.push('shader-f16'); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kArgumentTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable subgroups_f16;\nenable f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = quadBroadcast(${type.create(0).wgsl()}, 0); +}`; + + t.expectCompileResult(elementTypeOf(type) !== Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates data parameter type') + .params(u => + u + .combine('dataType', keysOf(kArgumentTypes)) + .combine('retType', keysOf(kArgumentTypes)) + .filter(t => { + const retType = kArgumentTypes[t.retType]; + const retEleTy = elementTypeOf(retType); + const dataType = kArgumentTypes[t.dataType]; + const dataEleTy = elementTypeOf(dataType); + return ( + retEleTy !== Type.abstractInt && + retEleTy !== Type.abstractFloat && + dataEleTy !== Type.abstractInt && + dataEleTy !== Type.abstractFloat + ); + }) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const dataType = kArgumentTypes[t.params.dataType]; + const retType = kArgumentTypes[t.params.retType]; + if (dataType.requiresF16() || retType.requiresF16()) { + features.push('subgroups-f16' as GPUFeatureName); + features.push('shader-f16'); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const dataType = kArgumentTypes[t.params.dataType]; + const retType = kArgumentTypes[t.params.retType]; + let enables = `enable subgroups;\n`; + if (dataType.requiresF16() || retType.requiresF16()) { + enables += `enable subgroups_f16;\nenable f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = quadBroadcast(${dataType.create(0).wgsl()}, 0); +}`; + + const expect = elementTypeOf(dataType) !== Type.bool && dataType === retType; + t.expectCompileResult(expect, wgsl); + }); + +g.test('id_type') + .desc('Validates id parameter type') + .params(u => u.combine('type', keysOf(kArgumentTypes))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const type = kArgumentTypes[t.params.type]; + const wgsl = ` +enable subgroups; +@compute @workgroup_size(1) +fn main() { + _ = quadBroadcast(0, ${type.create(0).wgsl()}); +}`; + + const expect = isConvertible(type, Type.u32) || isConvertible(type, Type.i32); + t.expectCompileResult(expect, wgsl); + }); + +const kIdCases = { + const_decl: { + code: 'const_decl', + valid: true, + }, + const_literal: { + code: '0', + valid: true, + }, + const_expr: { + code: 'const_decl + 2', + valid: true, + }, + let_decl: { + code: 'let_decl', + valid: false, + }, + override_decl: { + code: 'override_decl', + valid: false, + }, + var_func_decl: { + code: 'var_func_decl', + valid: false, + }, + var_priv_decl: { + code: 'var_priv_decl', + valid: false, + }, +}; + +g.test('id_constness') + .desc('Validates that id must be a const-expression') + .params(u => u.combine('value', keysOf(kIdCases))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +override override_decl : u32; +var var_priv_decl : u32; +fn foo() { + var var_func_decl : u32; + let let_decl = var_func_decl; + const const_decl = 0u; + _ = quadBroadcast(0, ${kIdCases[t.params.value].code}); +}`; + + t.expectCompileResult(kIdCases[t.params.value].valid, wgsl); + }); + +g.test('stage') + .desc('Validates it is only usable in correct stage') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = quadBroadcast(0, 0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/quadSwap.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/quadSwap.spec.ts new file mode 100644 index 000000000000..3812ba057ed6 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/quadSwap.spec.ts @@ -0,0 +1,227 @@ +export const description = ` +Validation tests for quadSwapX, quadSwapY, and quadSwapDiagonal. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { + Type, + elementTypeOf, + kAllScalarsAndVectors, + isConvertible, +} from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kOps = ['quadSwapX', 'quadSwapY', 'quadSwapDiagonal'] as const; + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = ${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = ${t.params.op}(0h); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record string> = { + constant: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = ${op}(0); +}`; + }, + override: (op: string) => { + return ` +enable subgroups +override o = ${op}(0);`; + }, + runtime: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = ${op}(0); +}`; + }, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages)).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage](t.params.op); + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(${type.create(0).wgsl()}); +}`; + + const eleType = elementTypeOf(type); + t.expectCompileResult(eleType !== Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('retType', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.retType]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + .combine('op', kOps) + .combine('paramType', keysOf(kTypes)) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + if (retType.requiresF16() || paramType.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + let enables = `enable subgroups;\n`; + if (retType.requiresF16() || paramType.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = ${t.params.op}(${paramType.create(0).wgsl()}); +}`; + + // Can't just use isConvertible since functions must concretize the parameter + // type before examining the whole statement. + const eleParamType = elementTypeOf(paramType); + const eleRetType = elementTypeOf(retType); + let expect = paramType === retType && eleRetType !== Type.bool; + if (eleParamType === Type.abstractInt) { + expect = eleRetType === Type.i32 && isConvertible(paramType, retType); + } else if (eleParamType === Type.abstractFloat) { + expect = eleRetType === Type.f32 && isConvertible(paramType, retType); + } + t.expectCompileResult(expect, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = ${t.params.op}(0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupAnyAll.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupAnyAll.spec.ts new file mode 100644 index 000000000000..eaee33e62cff --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupAnyAll.spec.ts @@ -0,0 +1,186 @@ +export const description = ` +Validation tests for subgroupAny and subgroupAll. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kOps = ['subgroupAny', 'subgroupAll'] as const; + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = ${t.params.op}(true); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record string> = { + constant: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = ${op}(true); +}`; + }, + override: (op: string) => { + return ` +enable subgroups +override o = select(0, 1, ${op}(true));`; + }, + runtime: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = ${op}(true); +}`; + }, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages)).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage](t.params.op); + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}${t.params.op}(false); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(${type.create(0).wgsl()}); +}`; + + t.expectCompileResult(type === Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('type', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.type]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + .combine('op', kOps) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${type.toString()} = ${t.params.op}(true); +}`; + + t.expectCompileResult(type === Type.bool, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = ${t.params.op}(true); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupBallot.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupBallot.spec.ts index afbe33e93c56..5f53847be25c 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/subgroupBallot.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupBallot.spec.ts @@ -9,6 +9,22 @@ import { ShaderValidationTest } from '../../../shader_validation_test.js'; export const g = makeTestGroup(ShaderValidationTest); +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = subgroupBallot(true); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + const kStages: Record = { constant: ` enable subgroups; @@ -38,6 +54,23 @@ g.test('early_eval') t.expectCompileResult(t.params.stage === 'runtime', code); }); +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}subgroupBallot(true); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + const kArgumentTypes = objectsToRecord(kAllScalarsAndVectors); g.test('data_type') @@ -69,7 +102,7 @@ fn main() { }); g.test('return_type') - .desc('Validates data parameter type') + .desc('Validates return type') .params(u => u.combine('type', keysOf(kArgumentTypes)).filter(t => { const type = kArgumentTypes[t.type]; diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupBitwise.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupBitwise.spec.ts new file mode 100644 index 000000000000..ca0dfb6fd719 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupBitwise.spec.ts @@ -0,0 +1,204 @@ +export const description = ` +Validation tests for subgroupAnd, subgroupOr, and subgroupXor. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { + Type, + elementTypeOf, + kAllScalarsAndVectors, + isConvertible, +} from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kOps = ['subgroupAnd', 'subgroupOr', 'subgroupXor'] as const; + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = ${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record string> = { + constant: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = ${op}(0); +}`; + }, + override: (op: string) => { + return ` +enable subgroups +override o = ${op}(0);`; + }, + runtime: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = ${op}(0); +}`; + }, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages)).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage](t.params.op); + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(${type.create(0).wgsl()}); +}`; + + const eleType = elementTypeOf(type); + const expect = isConvertible(eleType, Type.u32) || isConvertible(eleType, Type.i32); + t.expectCompileResult(expect, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('retType', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.retType]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + .combine('op', kOps) + .combine('paramType', keysOf(kTypes)) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + if (retType.requiresF16() || paramType.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + let enables = `enable subgroups;\n`; + if (retType.requiresF16() || paramType.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = ${t.params.op}(${paramType.create(0).wgsl()}); +}`; + + // Can't just use isConvertible since functions must concretize the parameter + // type before examining the whole statement. + const eleParamType = elementTypeOf(paramType); + const eleRetType = elementTypeOf(retType); + let expect = paramType === retType && (eleRetType === Type.i32 || eleRetType === Type.u32); + if (eleParamType === Type.abstractInt) { + expect = eleRetType === Type.i32 && isConvertible(paramType, retType); + } + t.expectCompileResult(expect, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = ${t.params.op}(0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcast.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcast.spec.ts index a71b145092c8..fd76cd419b7f 100644 --- a/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcast.spec.ts +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcast.spec.ts @@ -14,6 +14,44 @@ import { ShaderValidationTest } from '../../../shader_validation_test.js'; export const g = makeTestGroup(ShaderValidationTest); +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = subgroupBroadcast(0, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = subgroupBroadcast(0h, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + const kArgumentTypes = objectsToRecord(kAllScalarsAndVectors); const kStages: Record = { @@ -156,6 +194,58 @@ fn main() { t.expectCompileResult(expect, wgsl); }); +const kIdCases = { + const_decl: { + code: 'const_decl', + valid: true, + }, + const_literal: { + code: '0', + valid: true, + }, + const_expr: { + code: 'const_decl + 2', + valid: true, + }, + let_decl: { + code: 'let_decl', + valid: false, + }, + override_decl: { + code: 'override_decl', + valid: false, + }, + var_func_decl: { + code: 'var_func_decl', + valid: false, + }, + var_priv_decl: { + code: 'var_priv_decl', + valid: false, + }, +}; + +g.test('id_constness') + .desc('Validates that id must be a const-expression') + .params(u => u.combine('value', keysOf(kIdCases))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +override override_decl : u32; +var var_priv_decl : u32; +fn foo() { + var var_func_decl : u32; + let let_decl = var_func_decl; + const const_decl = 0u; + _ = subgroupBroadcast(0, ${kIdCases[t.params.value].code}); +}`; + + t.expectCompileResult(kIdCases[t.params.value].valid, wgsl); + }); + g.test('stage') .desc('Validates it is only usable in correct stage') .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const)) diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcastFirst.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcastFirst.spec.ts new file mode 100644 index 000000000000..4525b6b97ef8 --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupBroadcastFirst.spec.ts @@ -0,0 +1,210 @@ +export const description = ` +Validation tests for subgroupBroadcastFirst +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = subgroupBroadcastFirst(0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = subgroupBroadcastFirst(0h); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kArgumentTypes = objectsToRecord(kAllScalarsAndVectors); + +const kStages: Record = { + constant: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = subgroupBroadcastFirst(0); +}`, + override: ` +enable subgroups; +override o = subgroupBroadcastFirst(0);`, + runtime: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = subgroupBroadcastFirst(0); +}`, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage]; + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}subgroupBroadcastFirst(0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kArgumentTypes))) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kArgumentTypes[t.params.type]; + if (type.requiresF16()) { + features.push('subgroups-f16' as GPUFeatureName); + features.push('shader-f16'); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kArgumentTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable subgroups_f16;\nenable f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = subgroupBroadcastFirst(${type.create(0).wgsl()}); +}`; + + t.expectCompileResult(elementTypeOf(type) !== Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates data parameter type') + .params(u => + u + .combine('dataType', keysOf(kArgumentTypes)) + .combine('retType', keysOf(kArgumentTypes)) + .filter(t => { + const retType = kArgumentTypes[t.retType]; + const retEleTy = elementTypeOf(retType); + const dataType = kArgumentTypes[t.dataType]; + const dataEleTy = elementTypeOf(dataType); + return ( + retEleTy !== Type.abstractInt && + retEleTy !== Type.abstractFloat && + dataEleTy !== Type.abstractInt && + dataEleTy !== Type.abstractFloat + ); + }) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const dataType = kArgumentTypes[t.params.dataType]; + const retType = kArgumentTypes[t.params.retType]; + if (dataType.requiresF16() || retType.requiresF16()) { + features.push('subgroups-f16' as GPUFeatureName); + features.push('shader-f16'); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const dataType = kArgumentTypes[t.params.dataType]; + const retType = kArgumentTypes[t.params.retType]; + let enables = `enable subgroups;\n`; + if (dataType.requiresF16() || retType.requiresF16()) { + enables += `enable subgroups_f16;\nenable f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = subgroupBroadcastFirst(${dataType.create(0).wgsl()}); +}`; + + const expect = elementTypeOf(dataType) !== Type.bool && dataType === retType; + t.expectCompileResult(expect, wgsl); + }); + +g.test('stage') + .desc('Validates it is only usable in correct stage') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = subgroupBroadcastFirst(0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupElect.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupElect.spec.ts new file mode 100644 index 000000000000..5637860c59ce --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupElect.spec.ts @@ -0,0 +1,175 @@ +export const description = ` +Validation tests for subgroupElect. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { Type, elementTypeOf, kAllScalarsAndVectors } from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = subgroupElect(); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record = { + constant: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = subgroupElect(); +}`, + override: ` +enable subgroups +override o = select(0, 1, subgroupElect());`, + runtime: ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = subgroupElect(); +}`, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages))) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage]; + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}subgroupElect(); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates there are no valid data parameters') + .params(u => u.combine('type', keysOf(kTypes))) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = subgroupElect(${type.create(0).wgsl()}); +}`; + + t.expectCompileResult(false, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u.combine('type', keysOf(kTypes)).filter(t => { + const type = kTypes[t.type]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${type.toString()} = subgroupElect(); +}`; + + t.expectCompileResult(type === Type.bool, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = subgroupElect(); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupMinMax.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupMinMax.spec.ts new file mode 100644 index 000000000000..84c1860019ee --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupMinMax.spec.ts @@ -0,0 +1,227 @@ +export const description = ` +Validation tests for subgroupMin and subgroupMax. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { + Type, + elementTypeOf, + kAllScalarsAndVectors, + isConvertible, +} from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kOps = ['subgroupMin', 'subgroupMax'] as const; + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = ${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = ${t.params.op}(0h); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record string> = { + constant: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = ${op}(0); +}`; + }, + override: (op: string) => { + return ` +enable subgroups +override o = ${op}(0);`; + }, + runtime: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = ${op}(0); +}`; + }, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages)).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage](t.params.op); + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}${t.params.op}(0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(${type.create(0).wgsl()}); +}`; + + const eleType = elementTypeOf(type); + t.expectCompileResult(eleType !== Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('retType', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.retType]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + .combine('op', kOps) + .combine('paramType', keysOf(kTypes)) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + if (retType.requiresF16() || paramType.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + let enables = `enable subgroups;\n`; + if (retType.requiresF16() || paramType.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = ${t.params.op}(${paramType.create(0).wgsl()}); +}`; + + // Can't just use isConvertible since functions must concretize the parameter + // type before examining the whole statement. + const eleParamType = elementTypeOf(paramType); + const eleRetType = elementTypeOf(retType); + let expect = paramType === retType && eleRetType !== Type.bool; + if (eleParamType === Type.abstractInt) { + expect = eleRetType === Type.i32 && isConvertible(paramType, retType); + } else if (eleParamType === Type.abstractFloat) { + expect = eleRetType === Type.f32 && isConvertible(paramType, retType); + } + t.expectCompileResult(expect, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = ${t.params.op}(0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + }); diff --git a/src/webgpu/shader/validation/expression/call/builtin/subgroupShuffle.spec.ts b/src/webgpu/shader/validation/expression/call/builtin/subgroupShuffle.spec.ts new file mode 100644 index 000000000000..62ffb5af36dd --- /dev/null +++ b/src/webgpu/shader/validation/expression/call/builtin/subgroupShuffle.spec.ts @@ -0,0 +1,262 @@ +export const description = ` +Validation tests for subgroupShuffle, subgroupShuffleXor, subgroupShuffleUp, and subgroupShuffleDown. +`; + +import { makeTestGroup } from '../../../../../../common/framework/test_group.js'; +import { keysOf, objectsToRecord } from '../../../../../../common/util/data_tables.js'; +import { + Type, + elementTypeOf, + kAllScalarsAndVectors, + isConvertible, +} from '../../../../../util/conversion.js'; +import { ShaderValidationTest } from '../../../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kOps = [ + 'subgroupShuffle', + 'subgroupShuffleXor', + 'subgroupShuffleUp', + 'subgroupShuffleDown', +] as const; + +g.test('requires_subgroups') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +${t.params.enable ? 'enable subgroups;' : ''} +fn foo() { + _ = ${t.params.op}(0, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +g.test('requires_subgroups_f16') + .desc('Validates that the subgroups feature is required') + .params(u => u.combine('enable', [false, true] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + const features: GPUFeatureName[] = ['shader-f16', 'subgroups' as GPUFeatureName]; + if (t.params.enable) { + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const wgsl = ` +enable f16; +enable subgroups; +${t.params.enable ? 'enable subgroups_f16;' : ''} +fn foo() { + _ = ${t.params.op}(0h, 0); +}`; + + t.expectCompileResult(t.params.enable, wgsl); + }); + +const kStages: Record string> = { + constant: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + const x = ${op}(0, 0); +}`; + }, + override: (op: string) => { + return ` +enable subgroups +override o = ${op}(0, 0);`; + }, + runtime: (op: string) => { + return ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + let x = ${op}(0, 0); +}`; + }, +}; + +g.test('early_eval') + .desc('Ensures the builtin is not able to be compile time evaluated') + .params(u => u.combine('stage', keysOf(kStages)).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const code = kStages[t.params.stage](t.params.op); + t.expectCompileResult(t.params.stage === 'runtime', code); + }); + +g.test('must_use') + .desc('Tests that the builtin has the @must_use attribute') + .params(u => u.combine('must_use', [true, false] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const wgsl = ` +enable subgroups; +@compute @workgroup_size(16) +fn main() { + ${t.params.must_use ? '_ = ' : ''}${t.params.op}(0, 0); +}`; + + t.expectCompileResult(t.params.must_use, wgsl); + }); + +const kTypes = objectsToRecord(kAllScalarsAndVectors); + +g.test('data_type') + .desc('Validates data parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(${type.create(0).wgsl()}, 0); +}`; + + const eleType = elementTypeOf(type); + t.expectCompileResult(eleType !== Type.bool, wgsl); + }); + +g.test('return_type') + .desc('Validates return type') + .params(u => + u + .combine('retType', keysOf(kTypes)) + .filter(t => { + const type = kTypes[t.retType]; + const eleType = elementTypeOf(type); + return eleType !== Type.abstractInt && eleType !== Type.abstractFloat; + }) + .combine('op', kOps) + .combine('paramType', keysOf(kTypes)) + ) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + if (retType.requiresF16() || paramType.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const retType = kTypes[t.params.retType]; + const paramType = kTypes[t.params.paramType]; + let enables = `enable subgroups;\n`; + if (retType.requiresF16() || paramType.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + let res : ${retType.toString()} = ${t.params.op}(${paramType.create(0).wgsl()}, 0); +}`; + + // Can't just use isConvertible since functions must concretize the parameter + // type before examining the whole statement. + const eleParamType = elementTypeOf(paramType); + const eleRetType = elementTypeOf(retType); + let expect = paramType === retType && eleRetType !== Type.bool; + if (eleParamType === Type.abstractInt) { + expect = eleRetType === Type.i32 && isConvertible(paramType, retType); + } else if (eleParamType === Type.abstractFloat) { + expect = eleRetType === Type.f32 && isConvertible(paramType, retType); + } + t.expectCompileResult(expect, wgsl); + }); + +g.test('param2_type') + .desc('Validates shuffle parameter type') + .params(u => u.combine('type', keysOf(kTypes)).combine('op', kOps)) + .beforeAllSubcases(t => { + const features = ['subgroups' as GPUFeatureName]; + const type = kTypes[t.params.type]; + if (type.requiresF16()) { + features.push('shader-f16'); + features.push('subgroups-f16' as GPUFeatureName); + } + t.selectDeviceOrSkipTestCase(features); + }) + .fn(t => { + const type = kTypes[t.params.type]; + let enables = `enable subgroups;\n`; + if (type.requiresF16()) { + enables += `enable f16;\nenable subgroups_f16;`; + } + const wgsl = ` +${enables} +@compute @workgroup_size(1) +fn main() { + _ = ${t.params.op}(0, ${type.create(0).wgsl()}); +}`; + + const expect = + isConvertible(type, Type.u32) || (type === Type.i32 && t.params.op === 'subgroupShuffle'); + t.expectCompileResult(expect, wgsl); + }); + +g.test('stage') + .desc('validates builtin is only usable in the correct stages') + .params(u => u.combine('stage', ['compute', 'fragment', 'vertex'] as const).combine('op', kOps)) + .beforeAllSubcases(t => { + t.selectDeviceOrSkipTestCase('subgroups' as GPUFeatureName); + }) + .fn(t => { + const compute = ` +@compute @workgroup_size(1) +fn main() { + foo(); +}`; + + const fragment = ` +@fragment +fn main() { + foo(); +}`; + + const vertex = ` +@vertex +fn main() -> @builtin(position) vec4f { + foo(); + return vec4f(); +}`; + + const entry = { compute, fragment, vertex }[t.params.stage]; + const wgsl = ` +enable subgroups; +fn foo() { + _ = ${t.params.op}(0, 0); +} + +${entry} +`; + + t.expectCompileResult(t.params.stage !== 'vertex', wgsl); + });