Skip to content
Open
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
125 changes: 125 additions & 0 deletions src/webgpu/shader/validation/extension/subgroup_size_control.spec.ts
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
Copy link
Copy Markdown
Collaborator

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.

Copy link
Copy Markdown
Contributor

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.

* 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);
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The 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.

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The 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;' : ''}
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The 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
${enableExtension ? 'enable subgroups; enable subgroup_size_control;' : ''}
enable subgroups;
${enableExtension ? 'enable subgroup_size_control;' : ''}

@compute @workgroup_size(${subgroupSize}) @subgroup_size(${subgroupSize})
fn main() {}
`
);
});

const kStageShaders = {
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The 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));
});
Loading