export const description = `
createComputePipeline and createComputePipelineAsync validation tests.
Note: entry point matching tests are in shader_module/entry_point.spec.ts
`;
import { AllFeaturesMaxLimitsGPUTest } from '../.././gpu_test.js';
import { makeTestGroup } from '../../../common/framework/test_group.js';
import { keysOf } from '../../../common/util/data_tables.js';
import {
isTextureFormatUsableWithStorageAccessMode,
kPossibleStorageTextureFormats,
} from '../../format_info.js';
import { kValue } from '../../util/constants.js';
import { TShaderStage, getShaderWithEntryPoint } from '../../util/shader.js';
import {
kAPIResources,
getWGSLShaderForResource,
getAPIBindGroupLayoutForResource,
doResourcesMatch,
} from './utils.js';
import * as vtu from './validation_test_utils.js';
class F extends AllFeaturesMaxLimitsGPUTest {
getShaderModule(
shaderStage: TShaderStage = 'compute',
entryPoint: string = 'main'
): GPUShaderModule {
return this.device.createShaderModule({
code: getShaderWithEntryPoint(shaderStage, entryPoint),
});
}
}
export const g = makeTestGroup(F);
g.test('basic')
.desc(
`
Control case for createComputePipeline and createComputePipelineAsync.
Call the API with valid compute shader and matching valid entryPoint, making sure that the test function working well.
`
)
.params(u => u.combine('isAsync', [true, false]))
.fn(t => {
const { isAsync } = t.params;
vtu.doCreateComputePipelineTest(t, isAsync, true, {
layout: 'auto',
compute: { module: t.getShaderModule('compute', 'main'), entryPoint: 'main' },
});
});
g.test('shader_module,invalid')
.desc(
`
Tests calling createComputePipeline(Async) with a invalid compute shader, and check that the APIs catch this error.
`
)
.params(u => u.combine('isAsync', [true, false]))
.fn(t => {
const { isAsync } = t.params;
vtu.doCreateComputePipelineTest(t, isAsync, false, {
layout: 'auto',
compute: {
module: vtu.createInvalidShaderModule(t),
entryPoint: 'main',
},
});
});
g.test('shader_module,compute')
.desc(
`
Tests calling createComputePipeline(Async) with valid but different stage shader and matching entryPoint,
and check that the APIs only accept compute shader.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('shaderModuleStage', ['compute', 'vertex', 'fragment'] as TShaderStage[])
)
.fn(t => {
const { isAsync, shaderModuleStage } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.getShaderModule(shaderModuleStage, 'main'),
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(t, isAsync, shaderModuleStage === 'compute', descriptor);
});
g.test('shader_module,device_mismatch')
.desc(
'Tests createComputePipeline(Async) cannot be called with a shader module created from another device'
)
.paramsSubcasesOnly(u => u.combine('isAsync', [true, false]).combine('mismatched', [true, false]))
.beforeAllSubcases(t => t.usesMismatchedDevice())
.fn(t => {
const { isAsync, mismatched } = t.params;
const sourceDevice = mismatched ? t.mismatchedDevice : t.device;
const module = sourceDevice.createShaderModule({
code: '@compute @workgroup_size(1) fn main() {}',
});
const descriptor = {
layout: 'auto' as const,
compute: {
module,
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(t, isAsync, !mismatched, descriptor);
});
g.test('pipeline_layout,device_mismatch')
.desc(
'Tests createComputePipeline(Async) cannot be called with a pipeline layout created from another device'
)
.paramsSubcasesOnly(u => u.combine('isAsync', [true, false]).combine('mismatched', [true, false]))
.beforeAllSubcases(t => t.usesMismatchedDevice())
.fn(t => {
const { isAsync, mismatched } = t.params;
const sourceDevice = mismatched ? t.mismatchedDevice : t.device;
const layout = sourceDevice.createPipelineLayout({ bindGroupLayouts: [] });
const descriptor = {
layout,
compute: {
module: t.getShaderModule('compute', 'main'),
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(t, isAsync, !mismatched, descriptor);
});
g.test('limits,workgroup_storage_size')
.desc(
`
Tests calling createComputePipeline(Async) validation for compute using <= device.limits.maxComputeWorkgroupStorageSize bytes of workgroup storage.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ type: 'vec4<f32>', _typeSize: 16 },
{ type: 'mat4x4<f32>', _typeSize: 64 },
])
.beginSubcases()
.combine('countDeltaFromLimit', [0, 1])
)
.fn(t => {
const { isAsync, type, _typeSize, countDeltaFromLimit } = t.params;
const countAtLimit = Math.floor(t.device.limits.maxComputeWorkgroupStorageSize / _typeSize);
const count = countAtLimit + countDeltaFromLimit;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
var<workgroup> data: array<${type}, ${count}>;
@compute @workgroup_size(64) fn main () {
_ = data;
}
`,
}),
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(t, isAsync, count <= countAtLimit, descriptor);
});
g.test('limits,invocations_per_workgroup')
.desc(
`
Tests calling createComputePipeline(Async) validation for compute using <= device.limits.maxComputeInvocationsPerWorkgroup per workgroup.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('size', [
[128, 1, 2],
[129, 1, 2],
[2, 128, 1],
[2, 129, 1],
[1, 8, 32],
[1, 8, 33],
])
)
.fn(t => {
const { isAsync, size } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
@compute @workgroup_size(${size.join(',')}) fn main () {
}
`,
}),
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(
t,
isAsync,
size[0] * size[1] * size[2] <= t.device.limits.maxComputeInvocationsPerWorkgroup,
descriptor
);
});
g.test('limits,invocations_per_workgroup,each_component')
.desc(
`
Tests calling createComputePipeline(Async) validation for compute workgroup_size attribute has each component no more than their limits.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('size', [
[64],
[256, 1, 1],
[257, 1, 1],
[1, 256, 1],
[1, 257, 1],
[1, 1, 63],
[1, 1, 64],
[1, 1, 65],
] as const)
)
.fn(t => {
const { isAsync, size } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
@compute @workgroup_size(${size.join(',')}) fn main () {
}
`,
}),
entryPoint: 'main',
},
};
const workgroupX = size[0];
const workgroupY = size[1] ?? 1;
const workgroupZ = size[2] ?? 1;
const _success =
workgroupX <= t.device.limits.maxComputeWorkgroupSizeX &&
workgroupY <= t.device.limits.maxComputeWorkgroupSizeY &&
workgroupZ <= t.device.limits.maxComputeWorkgroupSizeZ;
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
g.test('overrides,identifier')
.desc(
`
Tests calling createComputePipeline(Async) validation for overridable constants identifiers.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ constants: {}, _success: true },
{ constants: { c0: 0 }, _success: true },
{ constants: { c0: 0, c1: 1 }, _success: true },
{ constants: { 'c0\0': 0 }, _success: false },
{ constants: { c9: 0 }, _success: false },
{ constants: { 1: 0 }, _success: true },
{ constants: { c3: 0 }, _success: false },
{ constants: { 2: 0 }, _success: false },
{ constants: { 1000: 0 }, _success: true },
{ constants: { 9999: 0 }, _success: false },
{ constants: { 1000: 0, c2: 0 }, _success: false },
{ constants: { 数: 0 }, _success: true },
{ constants: { séquençage: 0 }, _success: false },
] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
)
.fn(t => {
const { isAsync, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override c0: bool = true; // type: bool
override c1: u32 = 0u; // default override
override 数: u32 = 0u; // non-ASCII
override séquençage: u32 = 0u; // normalizable unicode (WGSL does not normalize)
@id(1000) override c2: u32 = 10u; // default
@id(1) override c3: u32 = 11u; // default
@compute @workgroup_size(1) fn main () {
// make sure the overridable constants are not optimized out
_ = u32(c0);
_ = u32(c1);
_ = u32(c2 + séquençage);
_ = u32(c3 + 数);
}`,
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
g.test('overrides,uninitialized')
.desc(
`
Tests calling createComputePipeline(Async) validation for uninitialized overridable constants.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ constants: {}, _success: false },
{ constants: { c0: 0, c2: 0, c8: 0 }, _success: false },
{ constants: { c0: 0, c2: 0, c5: 0, c8: 0 }, _success: true },
{ constants: { c0: 0, c2: 0, c5: 0, c8: 0, c1: 0 }, _success: true },
] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
)
.fn(t => {
const { isAsync, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override c0: bool; // type: bool
override c1: bool = false; // default override
override c2: f32; // type: float32
override c3: f32 = 0.0; // default override
override c4: f32 = 4.0; // default
override c5: i32; // type: int32
override c6: i32 = 0; // default override
override c7: i32 = 7; // default
override c8: u32; // type: uint32
override c9: u32 = 0u; // default override
@id(1000) override c10: u32 = 10u; // default
@compute @workgroup_size(1) fn main () {
// make sure the overridable constants are not optimized out
_ = u32(c0);
_ = u32(c1);
_ = u32(c2);
_ = u32(c3);
_ = u32(c4);
_ = u32(c5);
_ = u32(c6);
_ = u32(c7);
_ = u32(c8);
_ = u32(c9);
_ = u32(c10);
}`,
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
g.test('overrides,value,type_error')
.desc(
`
Tests calling createComputePipeline(Async) validation for constant values like inf, NaN will results in TypeError.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ constants: { cf: 1 }, _success: true },
{ constants: { cf: NaN }, _success: false },
{ constants: { cf: Number.POSITIVE_INFINITY }, _success: false },
{ constants: { cf: Number.NEGATIVE_INFINITY }, _success: false },
] as const)
)
.fn(t => {
const { isAsync, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override cf: f32 = 0.0;
@compute @workgroup_size(1) fn main () {
_ = cf;
}`,
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor, 'TypeError');
});
g.test('overrides,value,validation_error')
.desc(
`
Tests calling createComputePipeline(Async) validation for unrepresentable constant values in compute stage.
TODO(#2060): test with last_castable_pipeline_override.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ constants: { cu: kValue.u32.min }, _success: true },
{ constants: { cu: kValue.u32.min - 1 }, _success: false },
{ constants: { cu: kValue.u32.max }, _success: true },
{ constants: { cu: kValue.u32.max + 1 }, _success: false },
{ constants: { ci: kValue.i32.negative.min }, _success: true },
{ constants: { ci: kValue.i32.negative.min - 1 }, _success: false },
{ constants: { ci: kValue.i32.positive.max }, _success: true },
{ constants: { ci: kValue.i32.positive.max + 1 }, _success: false },
{ constants: { cf: kValue.f32.negative.min }, _success: true },
{
constants: { cf: kValue.f32.negative.first_non_castable_pipeline_override },
_success: false,
},
{ constants: { cf: kValue.f32.positive.max }, _success: true },
{
constants: { cf: kValue.f32.positive.first_non_castable_pipeline_override },
_success: false,
},
{ constants: { cb: Number.MAX_VALUE }, _success: true },
{ constants: { cb: kValue.i32.negative.min - 1 }, _success: true },
] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
)
.fn(t => {
const { isAsync, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override cb: bool = false;
override cu: u32 = 0u;
override ci: i32 = 0;
override cf: f32 = 0.0;
@compute @workgroup_size(1) fn main () {
_ = cb;
_ = cu;
_ = ci;
_ = cf;
}`,
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
g.test('overrides,entry_point,validation_error')
.desc(
`
Tests that pipeline constant (override) errors only trigger on entry point usage.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('pipeEntryPoint', ['main_success', 'main_pipe_error'])
)
.fn(t => {
const { isAsync, pipeEntryPoint } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override cu: u32 = 0u;
override cx: u32 = 1u/cu;
@compute @workgroup_size(1) fn main_success () {
_ = cu;
}
@compute @workgroup_size(1) fn main_pipe_error () {
_ = cx;
}`,
}),
entryPoint: pipeEntryPoint,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, pipeEntryPoint === 'main_success', descriptor);
});
g.test('overrides,value,validation_error,f16')
.desc(
`
Tests calling createComputePipeline(Async) validation for unrepresentable f16 constant values in compute stage.
TODO(#2060): Tighten the cases around the valid/invalid boundary once we have WGSL spec
clarity on whether values like f16.positive.last_castable_pipeline_override would be valid. See issue.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combineWithParams([
{ constants: { cf16: kValue.f16.negative.min }, _success: true },
{
constants: { cf16: kValue.f16.negative.first_non_castable_pipeline_override },
_success: false,
},
{ constants: { cf16: kValue.f16.positive.max }, _success: true },
{
constants: { cf16: kValue.f16.positive.first_non_castable_pipeline_override },
_success: false,
},
{ constants: { cf16: kValue.f32.negative.min }, _success: false },
{ constants: { cf16: kValue.f32.positive.max }, _success: false },
{
constants: { cf16: kValue.f32.negative.first_non_castable_pipeline_override },
_success: false,
},
{
constants: { cf16: kValue.f32.positive.first_non_castable_pipeline_override },
_success: false,
},
] as const)
)
.fn(t => {
t.skipIfDeviceDoesNotHaveFeature('shader-f16');
const { isAsync, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
enable f16;
override cf16: f16 = 0.0h;
@compute @workgroup_size(1) fn main () {
_ = cf16;
}`,
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
const kOverridesWorkgroupSizeShaders = {
u32: `
override x: u32 = 1u;
override y: u32 = 1u;
override z: u32 = 1u;
@compute @workgroup_size(x, y, z) fn main () {
_ = 0u;
}
`,
i32: `
override x: i32 = 1;
override y: i32 = 1;
override z: i32 = 1;
@compute @workgroup_size(x, y, z) fn main () {
_ = 0u;
}
`,
};
g.test('overrides,workgroup_size')
.desc(
`
Tests calling createComputePipeline(Async) validation for overridable constants used for workgroup size.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('type', ['u32', 'i32'] as const)
.combineWithParams([
{ constants: {}, _success: true },
{ constants: { x: 0, y: 0, z: 0 }, _success: false },
{ constants: { x: 1, y: -1, z: 1 }, _success: false },
{ constants: { x: 1, y: 0, z: 0 }, _success: false },
{ constants: { x: 16, y: 1, z: 1 }, _success: true },
] as { constants: Record<string, GPUPipelineConstantValue>; _success: boolean }[])
)
.fn(t => {
const { isAsync, type, constants, _success } = t.params;
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: kOverridesWorkgroupSizeShaders[type],
}),
entryPoint: 'main',
constants,
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
});
g.test('overrides,workgroup_size,limits')
.desc(
`
Tests calling createComputePipeline(Async) validation for overridable constants for workgroupSize exceeds device limits.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
.combine('type', ['u32', 'i32'] as const)
)
.fn(t => {
const { isAsync, type } = t.params;
const limits = t.device.limits;
const testFn = (x: number, y: number, z: number, _success: boolean) => {
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: kOverridesWorkgroupSizeShaders[type],
}),
entryPoint: 'main',
constants: {
x,
y,
z,
},
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
};
testFn(limits.maxComputeWorkgroupSizeX, 1, 1, true);
testFn(limits.maxComputeWorkgroupSizeX + 1, 1, 1, false);
testFn(1, limits.maxComputeWorkgroupSizeY, 1, true);
testFn(1, limits.maxComputeWorkgroupSizeY + 1, 1, false);
testFn(1, 1, limits.maxComputeWorkgroupSizeZ, true);
testFn(1, 1, limits.maxComputeWorkgroupSizeZ + 1, false);
testFn(
limits.maxComputeWorkgroupSizeX,
limits.maxComputeWorkgroupSizeY,
limits.maxComputeWorkgroupSizeZ,
limits.maxComputeWorkgroupSizeX *
limits.maxComputeWorkgroupSizeY *
limits.maxComputeWorkgroupSizeZ <=
limits.maxComputeInvocationsPerWorkgroup
);
});
g.test('overrides,workgroup_size,limits,workgroup_storage_size')
.desc(
`
Tests calling createComputePipeline(Async) validation for overridable constants for workgroupStorageSize exceeds device limits.
`
)
.params(u =>
u
.combine('isAsync', [true, false])
)
.fn(t => {
const { isAsync } = t.params;
const limits = t.device.limits;
const kVec4Size = 16;
const maxVec4Count = limits.maxComputeWorkgroupStorageSize / kVec4Size;
const kMat4Size = 64;
const maxMat4Count = limits.maxComputeWorkgroupStorageSize / kMat4Size;
const testFn = (vec4Count: number, mat4Count: number, _success: boolean) => {
const descriptor = {
layout: 'auto' as const,
compute: {
module: t.device.createShaderModule({
code: `
override a: u32;
override b: u32;
${vec4Count <= 0 ? '' : 'var<workgroup> vec4_data: array<vec4<f32>, a>;'}
${mat4Count <= 0 ? '' : 'var<workgroup> mat4_data: array<mat4x4<f32>, b>;'}
@compute @workgroup_size(1) fn main() {
${vec4Count <= 0 ? '' : '_ = vec4_data[0];'}
${mat4Count <= 0 ? '' : '_ = mat4_data[0];'}
}`,
}),
entryPoint: 'main',
constants: {
a: vec4Count,
b: mat4Count,
},
},
};
vtu.doCreateComputePipelineTest(t, isAsync, _success, descriptor);
};
testFn(1, 1, true);
testFn(maxVec4Count + 1, 0, false);
testFn(0, maxMat4Count + 1, false);
});
g.test('resource_compatibility')
.desc(
'Tests validation of resource (bind group) compatibility between pipeline layout and WGSL shader'
)
.params(u =>
u
.combine('apiResource', keysOf(kAPIResources))
.beginSubcases()
.combine('isAsync', [true, false] as const)
.combine('wgslResource', keysOf(kAPIResources))
)
.fn(t => {
const apiResource = kAPIResources[t.params.apiResource];
const wgslResource = kAPIResources[t.params.wgslResource];
t.skipIf(
wgslResource.storageTexture !== undefined &&
wgslResource.storageTexture.access !== 'write-only' &&
!t.hasLanguageFeature('readonly_and_readwrite_storage_textures'),
'Storage textures require language feature'
);
t.skipIfTextureViewDimensionNotSupported(wgslResource.texture?.viewDimension);
const layout = t.device.createPipelineLayout({
bindGroupLayouts: [
getAPIBindGroupLayoutForResource(t.device, GPUShaderStage.COMPUTE, apiResource),
],
});
const descriptor = {
layout,
compute: {
module: t.device.createShaderModule({
code: getWGSLShaderForResource('compute', wgslResource),
}),
entryPoint: 'main',
},
};
vtu.doCreateComputePipelineTest(
t,
t.params.isAsync,
doResourcesMatch(apiResource, wgslResource),
descriptor
);
});
g.test('storage_texture,format')
.desc(
`
Test that a pipeline with auto layout and storage texture access combo that is not supported
generates a validation error at createComputePipeline(Async)
`
)
.params(u =>
u
.combine('format', kPossibleStorageTextureFormats)
.beginSubcases()
.combine('isAsync', [true, false] as const)
.combine('access', ['read', 'write', 'read_write'] as const)
.combine('dimension', ['1d', '2d', '3d'] as const)
)
.fn(t => {
const { format, isAsync, access, dimension } = t.params;
t.skipIfTextureFormatNotSupported(format);
const code = `
@group(0) @binding(0) var tex: texture_storage_${dimension}<${format}, ${access}>;
@compute @workgroup_size(1) fn main() {
_ = tex;
}
`;
const module = t.device.createShaderModule({ code });
const success = isTextureFormatUsableWithStorageAccessMode(t.device.features, format, access);
const descriptor: GPUComputePipelineDescriptor = {
layout: 'auto',
compute: { module },
};
vtu.doCreateComputePipelineTest(t, isAsync, success, descriptor);
});
