-
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.
Contributor
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. Agreed, I don't see anything in the spec PR that suggests that the devices constraints would cause a shader-creation error, just the power-of-two and multiple-of-workgroup-x-dim things which are not dependent on device. So in theory you could just pick something like 32 here for the purposes of these tests. |
||||||||
| 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.
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.
I believe it is true that on these devices, the subgroup size is 8 for fragment shaders but that is never a valid subgroup size for a compute shader. This is the reason we decided to say that it would be a spurious failure if you requested a subgroup size that is not supported (otherwise we have to add a second set of limits as previously proposed, which we decided against).
So no, unfortunately not all sizes between min and max will work even for trivial compute pipelines.