From 58d4b115479ef11fbf2e8d08372f953e9beb2665 Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Mon, 16 Feb 2026 17:16:20 -0500 Subject: [PATCH 1/2] CTS for linear_indexing language feature * Execution tests for new builtins * Validation tests for new builtins --- src/webgpu/capability_info.ts | 1 + .../shader_io/compute_builtins.spec.ts | 69 ++++++++++++++++--- .../validation/shader_io/builtins.spec.ts | 14 ++++ 3 files changed, 74 insertions(+), 10 deletions(-) diff --git a/src/webgpu/capability_info.ts b/src/webgpu/capability_info.ts index 297fa55ccbe2..e2cf0e71400d 100644 --- a/src/webgpu/capability_info.ts +++ b/src/webgpu/capability_info.ts @@ -981,6 +981,7 @@ export const kKnownWGSLLanguageFeatures = [ 'subgroup_id', 'subgroup_uniformity', 'swizzle_assignment', + 'linear_indexing', ] as const; export type WGSLLanguageFeature = (typeof kKnownWGSLLanguageFeatures)[number]; diff --git a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts index 2d57a0f4fb1f..08183a3626c2 100644 --- a/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts +++ b/src/webgpu/shader/execution/shader_io/compute_builtins.spec.ts @@ -34,6 +34,7 @@ g.test('inputs') .beginSubcases() ) .fn(t => { + const linear_indexing = t.hasLanguageFeature('linear_indexing'); const invocationsPerGroup = t.params.groupSize.x * t.params.groupSize.y * t.params.groupSize.z; const totalInvocations = invocationsPerGroup * t.params.numGroups.x * t.params.numGroups.y * t.params.numGroups.z; @@ -46,6 +47,8 @@ g.test('inputs') let global_id = ''; let group_id = ''; let num_groups = ''; + let global_index = ''; + let group_index = ''; switch (t.params.method) { case 'param': params = ` @@ -54,12 +57,18 @@ g.test('inputs') @builtin(global_invocation_id) global_id : vec3, @builtin(workgroup_id) group_id : vec3, @builtin(num_workgroups) num_groups : vec3, + ${linear_indexing ? '@builtin(global_invocation_index) global_index : u32,' : ''} + ${linear_indexing ? '@builtin(workgroup_index) group_index : u32,' : ''} `; local_id = 'local_id'; local_index = 'local_index'; global_id = 'global_id'; group_id = 'group_id'; num_groups = 'num_groups'; + if (linear_indexing) { + global_index = 'global_index'; + group_index = 'group_index'; + } break; case 'struct': structures = `struct Inputs { @@ -68,6 +77,8 @@ g.test('inputs') @builtin(global_invocation_id) global_id : vec3, @builtin(workgroup_id) group_id : vec3, @builtin(num_workgroups) num_groups : vec3, + ${linear_indexing ? '@builtin(global_invocation_index) global_index : u32,' : ''} + ${linear_indexing ? '@builtin(workgroup_index) group_index : u32,' : ''} };`; params = `inputs : Inputs`; local_id = 'inputs.local_id'; @@ -75,6 +86,10 @@ g.test('inputs') global_id = 'inputs.global_id'; group_id = 'inputs.group_id'; num_groups = 'inputs.num_groups'; + if (linear_indexing) { + global_index = 'inputs.global_index'; + group_index = 'inputs.group_index'; + } break; case 'mixed': structures = `struct InputsA { @@ -87,12 +102,19 @@ g.test('inputs') params = `@builtin(local_invocation_id) local_id : vec3, inputsA : InputsA, inputsB : InputsB, - @builtin(num_workgroups) num_groups : vec3,`; + @builtin(num_workgroups) num_groups : vec3, + ${linear_indexing ? '@builtin(global_invocation_index) global_index : u32,' : ''} + ${linear_indexing ? '@builtin(workgroup_index) group_index : u32,' : ''} + `; local_id = 'local_id'; local_index = 'inputsA.local_index'; global_id = 'inputsA.global_id'; group_id = 'inputsB.group_id'; num_groups = 'num_groups'; + if (linear_indexing) { + global_index = 'global_index'; + group_index = 'group_index'; + } break; } @@ -104,6 +126,8 @@ g.test('inputs') global_id: vec3u, group_id: vec3u, num_groups: vec3u, + ${linear_indexing ? 'global_index : u32,' : ''} + ${linear_indexing ? 'group_index : u32,' : ''} }; @group(0) @binding(0) var outputs : array; @@ -117,15 +141,17 @@ g.test('inputs') fn main( ${params} ) { - let group_index = ((${group_id}.z * ${num_groups}.y) + ${group_id}.y) * ${num_groups}.x + ${group_id}.x; - let global_index = group_index * ${invocationsPerGroup}u + ${local_index}; + let o_group_index = ((${group_id}.z * ${num_groups}.y) + ${group_id}.y) * ${num_groups}.x + ${group_id}.x; + let o_global_index = o_group_index * ${invocationsPerGroup}u + ${local_index}; var o: Outputs; o.local_id = ${local_id}; o.local_index = ${local_index}; o.global_id = ${global_id}; o.group_id = ${group_id}; o.num_groups = ${num_groups}; - outputs[global_index] = o; + ${linear_indexing ? `o.global_index = ${global_index};` : ``} + ${linear_indexing ? `o.group_index = ${group_index};` : ``} + outputs[o_global_index] = o; } `; @@ -145,7 +171,9 @@ g.test('inputs') const kGlobalIdOffset = 4; const kGroupIdOffset = 8; const kNumGroupsOffset = 12; - const kOutputElementSize = 16; + const kGlobalIndexOffset = 15; + const kGroupIndexOffset = 16; + const kOutputElementSize = linear_indexing ? 20 : 16; // Create the output buffers. const outputBuffer = t.createBufferTracked({ @@ -203,6 +231,21 @@ g.test('inputs') const localIndex = (lz * t.params.groupSize.y + ly) * t.params.groupSize.x + lx; const globalIndex = groupIndex * invocationsPerGroup + localIndex; const globalOffset = globalIndex * kOutputElementSize; + const gidX = gx * t.params.groupSize.x + lx; + const gidY = gy * t.params.groupSize.y + ly; + const gidZ = gz * t.params.groupSize.z + lz; + const globalLinearIndex = + gidX + + gidY * t.params.groupSize.x * t.params.numGroups.x + + gidZ * + t.params.groupSize.x * + t.params.numGroups.x * + t.params.groupSize.y * + t.params.numGroups.y; + const groupLinearIndex = + gx + + gy * t.params.numGroups.x + + gz * t.params.numGroups.x * t.params.numGroups.y; const expectEqual = (name: string, expected: number, actual: number) => { if (actual !== expected) { @@ -226,17 +269,23 @@ g.test('inputs') const error = checkVec3Value('local_id', kLocalIdOffset, { x: lx, y: ly, z: lz }) || - checkVec3Value('global_id', kGlobalIdOffset, { - x: gx * t.params.groupSize.x + lx, - y: gy * t.params.groupSize.y + ly, - z: gz * t.params.groupSize.z + lz, - }) || + checkVec3Value('global_id', kGlobalIdOffset, { x: gidX, y: gidY, z: gidZ }) || checkVec3Value('group_id', kGroupIdOffset, { x: gx, y: gy, z: gz }) || checkVec3Value('num_groups', kNumGroupsOffset, t.params.numGroups) || expectEqual( 'local_index', localIndex, output[globalOffset + kLocalIndexOffset] + ) || + expectEqual( + 'global_index', + globalLinearIndex, + output[globalOffset + kGlobalIndexOffset] + ) || + expectEqual( + 'group_index', + groupLinearIndex, + output[globalOffset + kGroupIndexOffset] ); if (error) { return error; diff --git a/src/webgpu/shader/validation/shader_io/builtins.spec.ts b/src/webgpu/shader/validation/shader_io/builtins.spec.ts index e7d5b1070823..1c4e11aafb0e 100644 --- a/src/webgpu/shader/validation/shader_io/builtins.spec.ts +++ b/src/webgpu/shader/validation/shader_io/builtins.spec.ts @@ -119,6 +119,20 @@ export const kBuiltins: readonly Builtin[] = [ enable: 'subgroups', requires: 'subgroup_id', }, + { + name: 'workgroup_index', + stage: 'compute', + io: 'in', + type: 'u32', + requires: 'linear_indexing', + }, + { + name: 'global_invocation_index', + stage: 'compute', + io: 'in', + type: 'u32', + requires: 'linear_indexing', + }, ] as const; // List of types to test against. From 46e93f07e74375a32baf4e219062536d9b3ddd7b Mon Sep 17 00:00:00 2001 From: Alan Baker Date: Sat, 21 Feb 2026 21:57:13 -0500 Subject: [PATCH 2/2] Add dispatch validation tests --- src/webgpu/api/validation/dispatch.spec.ts | 143 +++++++++++++++++++++ src/webgpu/listing_meta.json | 2 + 2 files changed, 145 insertions(+) create mode 100644 src/webgpu/api/validation/dispatch.spec.ts diff --git a/src/webgpu/api/validation/dispatch.spec.ts b/src/webgpu/api/validation/dispatch.spec.ts new file mode 100644 index 000000000000..7ebc1d131482 --- /dev/null +++ b/src/webgpu/api/validation/dispatch.spec.ts @@ -0,0 +1,143 @@ +export const description = ` +Compute dispatch validation tests. +`; + +import { AllFeaturesMaxLimitsGPUTest } from '../.././gpu_test.js'; +import { makeTestGroup } from '../../../common/framework/test_group.js'; + +export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest); + +g.test('dispatch,linear_indexing_range') + .desc('Tests validation of total invocations for linear_indexing built-in values') + .params(u => + u + .combine('builtin', ['global_invocation_index', 'workgroup_index'] as const) + .beginSubcases() + .combine('size', ['max', 'valid'] as const) + ) + .fn(t => { + // Other builtins are not tested due to onerous runtimes. + t.skipIf(!t.hasLanguageFeature('linear_indexing'), 'Missing linear_indexing language feature'); + + // Spec limits: + // - maxComputeWorkgroupsPerDimension = 65535 + const { maxComputeWorkgroupsPerDimension } = t.device.limits; + const x = t.params.builtin === 'global_invocation_index' ? 2 : 1, + y = 1, + z = 1; + const wgSize = x * y * z; + const countX = maxComputeWorkgroupsPerDimension; + const countY = t.params.size === 'max' ? maxComputeWorkgroupsPerDimension : 1; + const countZ = t.params.builtin === 'workgroup_index' ? 2 : 1; + + const totalInvocations = wgSize * countX * countY * countZ; + t.skipIf(t.params.size === 'max' && totalInvocations <= 0xffffffff, 'Uninteresting test'); + + const code = ` +@compute @workgroup_size(${x}, ${y}, ${z}) +fn main(@builtin(${t.params.builtin}) input : u32) { + _ = input; +}`; + + const shaderModule = t.device.createShaderModule({ code }); + const computePipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: shaderModule, + }, + }); + const commandEncoder = t.device.createCommandEncoder(); + const computePassEncoder = commandEncoder.beginComputePass(); + computePassEncoder.setPipeline(computePipeline); + computePassEncoder.dispatchWorkgroups(countX, countY, countZ); + computePassEncoder.end(); + + t.expectValidationError(() => { + t.queue.submit([commandEncoder.finish()]); + }, t.params.size === 'max'); + }); + +g.test('dispatchIndirect,linear_indexing_range') + .desc('Tests dispatchIndirect skips when linear_indexing is out of range') + .params(u => + u + .combine('builtin', ['global_invocation_index', 'workgroup_index'] as const) + .beginSubcases() + .combine('size', ['max', 'valid'] as const) + ) + .fn(t => { + // Other builtins are not tested due to onerous runtimes. + t.skipIf(!t.hasLanguageFeature('linear_indexing'), 'Missing linear_indexing language feature'); + + // Spec limits: + // - maxComputeWorkgroupsPerDimension = 65535 + const { maxComputeWorkgroupsPerDimension } = t.device.limits; + const x = t.params.builtin === 'global_invocation_index' ? 2 : 1, + y = 1, + z = 1; + const wgSize = x * y * z; + const countX = maxComputeWorkgroupsPerDimension; + const countY = t.params.size === 'max' ? maxComputeWorkgroupsPerDimension : 1; + const countZ = t.params.builtin === 'workgroup_index' ? 2 : 1; + + const totalInvocations = wgSize * countX * countY * countZ; + t.skipIf(t.params.size === 'max' && totalInvocations <= 0xffffffff, 'Uninteresting test'); + + const kMagic = 0xdeadbeef; + const code = ` +@group(0) @binding(0) +var out : u32; + +@compute @workgroup_size(${x}, ${y}, ${z}) +fn main(@builtin(${t.params.builtin}) input : u32, + @builtin(global_invocation_id) gid : vec3u) { + _ = input; + if (gid.x == 0 && gid.y == 0 && gid.z == 0) { + out = ${kMagic}; + } +}`; + + const dispatchIndirectCounts = new Uint32Array(3); + dispatchIndirectCounts[0] = countX; + dispatchIndirectCounts[1] = countY; + dispatchIndirectCounts[2] = countZ; + const indirectBuffer = t.makeBufferWithContents( + dispatchIndirectCounts, + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.INDIRECT + ); + t.trackForCleanup(indirectBuffer); + const outputBuffer = t.makeBufferWithContents( + new Uint32Array([0]), + GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE + ); + t.trackForCleanup(outputBuffer); + + const shaderModule = t.device.createShaderModule({ code }); + const computePipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: shaderModule, + }, + }); + const bg = t.device.createBindGroup({ + layout: computePipeline.getBindGroupLayout(0), + entries: [ + { + binding: 0, + resource: { + buffer: outputBuffer, + }, + }, + ], + }); + const commandEncoder = t.device.createCommandEncoder(); + const computePassEncoder = commandEncoder.beginComputePass(); + computePassEncoder.setPipeline(computePipeline); + computePassEncoder.setBindGroup(0, bg); + computePassEncoder.dispatchWorkgroupsIndirect(indirectBuffer, 0); + computePassEncoder.end(); + t.queue.submit([commandEncoder.finish()]); + + const expected = t.params.size === 'max' ? 0 : kMagic; + t.expectGPUBufferValuesEqual(outputBuffer, new Uint32Array([expected])); + }); diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index acbdadfb759a..8ac343b4bfa4 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -435,6 +435,8 @@ "webgpu:api,validation,createView:texture_view_usage:*": { "subcaseMS": 3106.634 }, "webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,command_encoder:*": { "subcaseMS": 1.522 }, "webgpu:api,validation,debugMarker:push_pop_call_count_unbalance,render_compute_pass:*": { "subcaseMS": 0.601 }, + "webgpu:api,validation,dispatch:dispatch,linear_indexing_range:*": { "subcaseMS": 359.656 }, + "webgpu:api,validation,dispatch:dispatchIndirect,linear_indexing_range:*": { "subcaseMS": 320.426 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,invalid_query_set:*": { "subcaseMS": 0.201 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,query_index:*": { "subcaseMS": 0.201 }, "webgpu:api,validation,encoding,beginComputePass:timestampWrites,query_set_type:*": { "subcaseMS": 0.401 },