-
Notifications
You must be signed in to change notification settings - Fork 105
Add shader validation tests on subgroup-size-control - Part I
#4641
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
Changes from all commits
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change | ||||||
|---|---|---|---|---|---|---|---|---|
| @@ -0,0 +1,125 @@ | ||||||||
| export const description = ` | ||||||||
| Validation tests for the subgroup_size_control extension | ||||||||
| `; | ||||||||
|
|
||||||||
| import { makeTestGroup } from '../../../../common/framework/test_group.js'; | ||||||||
| import { UniqueFeaturesAndLimitsShaderValidationTest } from '../shader_validation_test.js'; | ||||||||
|
|
||||||||
| export const g = makeTestGroup(UniqueFeaturesAndLimitsShaderValidationTest); | ||||||||
|
|
||||||||
| /** | ||||||||
| * Returns a subgroup size value that is valid for use in the @subgroup_size | ||||||||
| * attribute on the current adapter. | ||||||||
| * | ||||||||
| * On Intel gen-12lp, subgroupMinSize may be 8 in fragment stages, which is below the allowed range | ||||||||
| * for `[WaveSize]` on D3D12 (can only be 16). subgroupMaxSize (16) is always within the explicit | ||||||||
| * range, so it is returned for that architecture. | ||||||||
| * On all other adapters, subgroupMinSize is returned as the conservative choice as on many D3D12 | ||||||||
| * drivers only `waveLaneCountMin` is reliable, while `waveLaneCountMax` is not. | ||||||||
| * | ||||||||
| * @param adapterInfo The GPUAdapterInfo of the current device's adapter. | ||||||||
| * @returns A power-of-two subgroup size valid for @subgroup_size on this adapter. | ||||||||
| */ | ||||||||
| export function getValidSubgroupSizeForSubgroupSizeAttribute(adapterInfo: GPUAdapterInfo): number { | ||||||||
| interface SubgroupAdapterInfo extends GPUAdapterInfo { | ||||||||
| subgroupMinSize: number; | ||||||||
| subgroupMaxSize: number; | ||||||||
| } | ||||||||
| const { vendor, architecture, subgroupMinSize, subgroupMaxSize } = | ||||||||
| adapterInfo as SubgroupAdapterInfo; | ||||||||
| return vendor === 'intel' && architecture === 'gen-12lp' ? subgroupMaxSize : subgroupMinSize; | ||||||||
| } | ||||||||
|
|
||||||||
| g.test('enable_subgroup_size_control_requires_subgroups') | ||||||||
| .desc( | ||||||||
| `Checks that enabling the WGSL extension subgroup_size_control without also enabling the | ||||||||
| subgroups extension is a compilation error.` | ||||||||
| ) | ||||||||
| .params(u => u.combine('enableSubgroups', [false, true] as const)) | ||||||||
| .beforeAllSubcases(t => { | ||||||||
| t.selectDeviceOrSkipTestCase({ | ||||||||
| requiredFeatures: ['subgroup-size-control' as GPUFeatureName], | ||||||||
| }); | ||||||||
| }) | ||||||||
| .fn(t => { | ||||||||
| const { enableSubgroups } = t.params; | ||||||||
|
|
||||||||
| t.expectCompileResult( | ||||||||
| enableSubgroups, | ||||||||
| ` | ||||||||
| ${enableSubgroups ? 'enable subgroups;' : ''} | ||||||||
| enable subgroup_size_control; | ||||||||
| @compute @workgroup_size(1) | ||||||||
| fn main() {} | ||||||||
| ` | ||||||||
| ); | ||||||||
| }); | ||||||||
|
|
||||||||
| g.test('use_subgroup_size_attribute_requires_subgroup_size_control_extension_enabled') | ||||||||
| .desc( | ||||||||
| `Checks that the @subgroup_size attribute is only allowed with the WGSL extension | ||||||||
| subgroup_size_control enabled in the shader and the WebGPU extension subgroup-size-control | ||||||||
| supported on the device.` | ||||||||
| ) | ||||||||
| .params(u => u.combine('enableExtension', [false, true] as const)) | ||||||||
| .beforeAllSubcases(t => { | ||||||||
| t.selectDeviceOrSkipTestCase({ | ||||||||
| requiredFeatures: ['subgroup-size-control' as GPUFeatureName], | ||||||||
| }); | ||||||||
| }) | ||||||||
| .fn(t => { | ||||||||
| const { enableExtension } = t.params; | ||||||||
|
|
||||||||
| const subgroupSize = getValidSubgroupSizeForSubgroupSizeAttribute(t.device.adapterInfo); | ||||||||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Does compiling a shader really require the subgroup size to be valid for the device? It seems like it should work for any size, and only be validated at pipeline creation. |
||||||||
| t.expectCompileResult( | ||||||||
| enableExtension, | ||||||||
| ` | ||||||||
| ${enableExtension ? 'enable subgroups; enable subgroup_size_control;' : ''} | ||||||||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Tighten this so we're changing as little as possible between the true and false cases.
Suggested change
|
||||||||
| @compute @workgroup_size(${subgroupSize}) @subgroup_size(${subgroupSize}) | ||||||||
| fn main() {} | ||||||||
| ` | ||||||||
| ); | ||||||||
| }); | ||||||||
|
|
||||||||
| const kStageShaders = { | ||||||||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. nit: If these are only going to be used by one test, put them inside the test. |
||||||||
| compute: (subgroupSize: number) => ` | ||||||||
| enable subgroups; | ||||||||
| enable subgroup_size_control; | ||||||||
| @compute @workgroup_size(${subgroupSize}) @subgroup_size(${subgroupSize}) | ||||||||
| fn main() {} | ||||||||
| `, | ||||||||
| vertex: (subgroupSize: number) => ` | ||||||||
| enable subgroups; | ||||||||
| enable subgroup_size_control; | ||||||||
| @vertex @subgroup_size(${subgroupSize}) | ||||||||
| fn main() -> @builtin(position) vec4f { | ||||||||
| return vec4f(0); | ||||||||
| } | ||||||||
| `, | ||||||||
| fragment: (subgroupSize: number) => ` | ||||||||
| enable subgroups; | ||||||||
| enable subgroup_size_control; | ||||||||
| @fragment @subgroup_size(${subgroupSize}) | ||||||||
| fn main() -> @location(0) vec4f { | ||||||||
| return vec4f(0); | ||||||||
| } | ||||||||
| `, | ||||||||
| } as const; | ||||||||
|
|
||||||||
| g.test('subgroup_size_attribute_only_valid_in_compute_stage') | ||||||||
| .desc( | ||||||||
| `Checks that the @subgroup_size attribute is only valid on a compute shader entry point. | ||||||||
| Applying it to a vertex or fragment entry point must be a compilation error.` | ||||||||
| ) | ||||||||
| .params(u => u.combine('stage', ['compute', 'vertex', 'fragment'] as const)) | ||||||||
| .beforeAllSubcases(t => { | ||||||||
| t.selectDeviceOrSkipTestCase({ | ||||||||
| requiredFeatures: ['subgroup-size-control' as GPUFeatureName], | ||||||||
| }); | ||||||||
| }) | ||||||||
| .fn(t => { | ||||||||
| const { stage } = t.params; | ||||||||
| const subgroupSize = getValidSubgroupSizeForSubgroupSizeAttribute(t.device.adapterInfo); | ||||||||
|
|
||||||||
| t.expectCompileResult(stage === 'compute', kStageShaders[stage](subgroupSize)); | ||||||||
| }); | ||||||||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
You're saying adapterInfo.subgroupMinSize is 8 on these devices, but 8 is not actually a valid subgroup size in compute? That's kind of weird, it requires applications to do weird stuff like this to get the right subgroup size. I would have thought subgroupMinSize/subgroupMaxSize are just for compute, especially now that we're adding subgroup-size-control which only applies to compute.
I understand we're saying not all values between subgroupMinSize and subgroupMaxSize are valid to create all pipelines with, but surely they should all work for a trivial compute pipeline? Maybe we can't technically specify that, but I think we should try to test it unless there are known counterexamples.