From 0b46218c1a6cc779eb1be8726d24b97705c1fa55 Mon Sep 17 00:00:00 2001 From: Ib Green Date: Tue, 4 Aug 2026 21:40:17 -0400 Subject: [PATCH] fix(experimental): scale visibility workflow dispatches --- .../src/gpu-primitives/gpu-compaction.ts | 100 +++++++++---- .../src/gpu-primitives/gpu-mask.ts | 73 ++++++--- .../gpu-primitives/gpu-visibility-workflow.ts | 128 +++++++++++----- .../gpu-visibility-workflow.node.spec.ts | 139 ++++++++++++++++++ .../gpu-visibility-workflow.spec.ts | 117 +++++++++++++++ 5 files changed, 470 insertions(+), 87 deletions(-) create mode 100644 modules/experimental/test/gpu-primitives/gpu-visibility-workflow.node.spec.ts diff --git a/modules/experimental/src/gpu-primitives/gpu-compaction.ts b/modules/experimental/src/gpu-primitives/gpu-compaction.ts index 4c9f87f5a5..65202792dc 100644 --- a/modules/experimental/src/gpu-primitives/gpu-compaction.ts +++ b/modules/experimental/src/gpu-primitives/gpu-compaction.ts @@ -4,7 +4,12 @@ import {Computation} from '@luma.gl/engine'; import {GPUCommandGraph, GraphVectorView, type GraphDataView} from './gpu-command-graph'; -import {GPUScan} from './gpu-scan'; +import { + getBoundedDispatchLayout, + getBoundedInvocationIndexSource, + type GPUBoundedDispatchLayout +} from './gpu-dispatch-utils'; +import {addGPUScanToGraphWithDispatchLimit, GPUScan} from './gpu-scan'; import { createTransientVectorView, createTransientView, @@ -98,29 +103,56 @@ export class GPUCompaction { * commands. */ addToGraph(graph: GPUCommandGraph): void { - for (const view of [ - ...getCompactionChunks(this.input), - ...getCompactionChunks(this.flags), - ...getCompactionChunks(this.output), - this.count - ]) { - if (view.buffer.graph !== graph) { - throw new Error(`${this.id} views must belong to the target graph`); - } - } + addGPUCompactionToGraphWithDispatchLimit( + this, + graph, + graph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} - if (this.input.length === 0) { - addClearCountPass(graph, this.id, this.count); - return; +/** Adds stable scan and scatter passes with an explicit dispatch limit. @internal */ +export function addGPUCompactionToGraphWithDispatchLimit( + compaction: GPUCompaction, + graph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + for (const view of [ + ...getCompactionChunks(compaction.input), + ...getCompactionChunks(compaction.flags), + ...getCompactionChunks(compaction.output), + compaction.count + ]) { + if (view.buffer.graph !== graph) { + throw new Error(`${compaction.id} views must belong to the target graph`); } + } - const offsets = - this.flags instanceof GraphVectorView - ? createTransientVectorView(graph, `${this.id}-offsets`, this.flags) - : createTransientView(graph, `${this.id}-offsets`, 'uint32', this.flags.length); - new GPUScan({id: `${this.id}-scan`, input: this.flags, output: offsets}).addToGraph(graph); - addScatterPasses(graph, this.id, this.input, this.flags, offsets, this.output, this.count); + if (compaction.input.length === 0) { + addClearCountPass(graph, compaction.id, compaction.count); + return; } + + const offsets = + compaction.flags instanceof GraphVectorView + ? createTransientVectorView(graph, `${compaction.id}-offsets`, compaction.flags) + : createTransientView(graph, `${compaction.id}-offsets`, 'uint32', compaction.flags.length); + const scan = new GPUScan({ + id: `${compaction.id}-scan`, + input: compaction.flags, + output: offsets + }); + addGPUScanToGraphWithDispatchLimit(scan, graph, maxComputeWorkgroupsPerDimension); + addScatterPasses( + graph, + compaction.id, + compaction.input, + compaction.flags, + offsets, + compaction.output, + compaction.count, + maxComputeWorkgroupsPerDimension + ); } /** Writes the required zero count for an empty input. */ @@ -162,7 +194,8 @@ function addScatterPasses( flags: GPUCompactionInput, offsets: GPUCompactionInput, output: GPUCompactionInput, - count: GraphDataView<'uint32'> + count: GraphDataView<'uint32'>, + maxComputeWorkgroupsPerDimension: number ): void { const inputChunks = getCompactionChunks(input); const flagChunks = getCompactionChunks(flags); @@ -196,7 +229,13 @@ function addScatterPasses( output: outputChunk, outputStart, outputEnd, - count: writesCount ? count : undefined + count: writesCount ? count : undefined, + dispatchLayout: getBoundedDispatchLayout( + 'GPUCompaction', + inputChunks[inputChunkIndex].length, + COMPACTION_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ) }); } } @@ -216,6 +255,7 @@ function addScatterPass( outputStart: number; outputEnd: number; count?: GraphDataView<'uint32'>; + dispatchLayout: GPUBoundedDispatchLayout; } ): void { const countBinding = props.count @@ -242,9 +282,10 @@ ${props.count ? `const COUNT_OFFSET: u32 = ${getViewElementOffset(props.count)}u ${countBinding} @compute @workgroup_size(${COMPACTION_WORKGROUP_SIZE}) fn main( - @builtin(global_invocation_id) globalId: vec3 + @builtin(local_invocation_index) localInvocationIndex: u32, + @builtin(workgroup_id) workgroupId: vec3 ) { - let index = globalId.x; + ${getBoundedInvocationIndexSource(props.dispatchLayout, COMPACTION_WORKGROUP_SIZE)} if (index >= ELEMENT_COUNT) { return; } let flag = min(flags[FLAGS_OFFSET + index], 1u); let outputIndex = offsets[OFFSETS_OFFSET + index]; @@ -270,7 +311,7 @@ ${countBinding} outputValues: props.output, ...(props.count ? {outputCount: props.count} : {}) }, - dispatchCount: Math.ceil(props.input.length / COMPACTION_WORKGROUP_SIZE) + dispatchLayout: props.dispatchLayout }); } @@ -285,7 +326,7 @@ function addCompactionPass( usage: 'storage-read' | 'storage-write' | 'storage-read-write'; }>; bindings: Record; - dispatchCount: number; + dispatchLayout: GPUBoundedDispatchLayout; } ): void { graph.addComputePass({ @@ -311,7 +352,12 @@ function addCompactionPass( bindings[name] = getViewBinding(view, getBuffer); } computation.setBindings(bindings); - computation.dispatch(computePass, props.dispatchCount); + computation.dispatch( + computePass, + props.dispatchLayout.x, + props.dispatchLayout.y, + props.dispatchLayout.z + ); }, destroy: () => computation.destroy() }; diff --git a/modules/experimental/src/gpu-primitives/gpu-mask.ts b/modules/experimental/src/gpu-primitives/gpu-mask.ts index c42744ef19..600d4d3dfa 100644 --- a/modules/experimental/src/gpu-primitives/gpu-mask.ts +++ b/modules/experimental/src/gpu-primitives/gpu-mask.ts @@ -10,6 +10,11 @@ import { type GraphBufferUse, type GraphDataView } from './gpu-command-graph'; +import { + getBoundedDispatchLayout, + getBoundedInvocationIndexSource, + type GPUBoundedDispatchLayout +} from './gpu-dispatch-utils'; import { getViewBinding, getViewElementOffset, @@ -97,26 +102,45 @@ export class GPUMask { * The caller remains responsible for graph compilation, command submission, and readback. */ addToGraph(graph: GPUCommandGraph): void { - const outputChunks = getMaskChunks(this.output); - const inputChunks = this.inputs.map(getMaskChunks); - for (const chunk of [...outputChunks, ...inputChunks.flat()]) { - if (chunk.buffer.graph !== graph) { - throw new Error(`${this.id} masks must belong to the target graph`); - } + addGPUMaskToGraphWithDispatchLimit( + this, + graph, + graph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} + +/** Adds source-aligned mask composition with an explicit dispatch limit. @internal */ +export function addGPUMaskToGraphWithDispatchLimit( + mask: GPUMask, + graph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + const outputChunks = getMaskChunks(mask.output); + const inputChunks = mask.inputs.map(getMaskChunks); + for (const chunk of [...outputChunks, ...inputChunks.flat()]) { + if (chunk.buffer.graph !== graph) { + throw new Error(`${mask.id} masks must belong to the target graph`); } + } - for (const [chunkIndex, output] of outputChunks.entries()) { - if (output.length === 0) { - continue; - } - const inputs = inputChunks.map(chunks => chunks[chunkIndex]); - addMaskPass(graph, { - id: this.output instanceof GraphVectorView ? `${this.id}-chunk-${chunkIndex}` : this.id, - inputs, - output, - operation: this.operation - }); + for (const [chunkIndex, output] of outputChunks.entries()) { + if (output.length === 0) { + continue; } + const inputs = inputChunks.map(chunks => chunks[chunkIndex]); + addMaskPass(graph, { + id: mask.output instanceof GraphVectorView ? `${mask.id}-chunk-${chunkIndex}` : mask.id, + inputs, + output, + operation: mask.operation, + dispatchLayout: getBoundedDispatchLayout( + 'GPUMask', + output.length, + MASK_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ) + }); } } @@ -133,6 +157,7 @@ function addMaskPass( inputs: readonly GraphDataView<'uint32'>[]; output: GraphDataView<'uint32'>; operation: GPUMaskOperation; + dispatchLayout: GPUBoundedDispatchLayout; } ): void { const inputDeclarations = props.inputs @@ -154,8 +179,11 @@ ${inputDeclarations} @group(0) @binding(${outputBinding}) var outputMask: array; @compute @workgroup_size(${MASK_WORKGROUP_SIZE}) -fn main(@builtin(global_invocation_id) globalId: vec3) { - let index = globalId.x; +fn main( + @builtin(local_invocation_index) localInvocationIndex: u32, + @builtin(workgroup_id) workgroupId: vec3 +) { + ${getBoundedInvocationIndexSource(props.dispatchLayout, MASK_WORKGROUP_SIZE)} if (index >= ELEMENT_COUNT) { return; } @@ -193,7 +221,12 @@ fn main(@builtin(global_invocation_id) globalId: vec3) { resolvedBindings[name] = getViewBinding(view, getBuffer); } computation.setBindings(resolvedBindings); - computation.dispatch(computePass, Math.ceil(props.output.length / MASK_WORKGROUP_SIZE)); + computation.dispatch( + computePass, + props.dispatchLayout.x, + props.dispatchLayout.y, + props.dispatchLayout.z + ); }, destroy: () => computation.destroy() }; diff --git a/modules/experimental/src/gpu-primitives/gpu-visibility-workflow.ts b/modules/experimental/src/gpu-primitives/gpu-visibility-workflow.ts index b7ea16a83b..a4c6d0ae32 100644 --- a/modules/experimental/src/gpu-primitives/gpu-visibility-workflow.ts +++ b/modules/experimental/src/gpu-primitives/gpu-visibility-workflow.ts @@ -5,8 +5,17 @@ import {type Binding} from '@luma.gl/core'; import {Computation} from '@luma.gl/engine'; import {GPUCommandGraph, GraphVectorView, type GraphDataView} from './gpu-command-graph'; -import {GPUCompaction, type GPUCompactionInput} from './gpu-compaction'; -import {GPUMask} from './gpu-mask'; +import { + addGPUCompactionToGraphWithDispatchLimit, + GPUCompaction, + type GPUCompactionInput +} from './gpu-compaction'; +import { + getBoundedDispatchLayout, + getBoundedInvocationIndexSource, + type GPUBoundedDispatchLayout +} from './gpu-dispatch-utils'; +import {addGPUMaskToGraphWithDispatchLimit, GPUMask} from './gpu-mask'; import { createTransientVectorView, createTransientView, @@ -125,43 +134,65 @@ export class GPUVisibilityWorkflow { * Adds mask composition, identity generation, scan, scatter, and count publication to a graph. */ addToGraph(graph: GPUCommandGraph): void { - const template = this.predicates[0].mask; - for (const view of [ - ...this.predicates.flatMap(predicate => getVisibilityChunks(predicate.mask)), - ...getVisibilityChunks(this.output), - ...(this.outputMask ? getVisibilityChunks(this.outputMask) : []), - ...(this.sourceIds ? getVisibilityChunks(this.sourceIds) : []), - this.count - ]) { - if (view.buffer.graph !== graph) { - throw new Error(`${this.id} views must belong to the target graph`); - } - } + addGPUVisibilityWorkflowToGraphWithDispatchLimit( + this, + graph, + graph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} - const finalMask = - this.outputMask ?? createTransientVisibilityInput(graph, `${this.id}-mask`, template); - if (finalMask !== template || this.predicates.length > 1) { - new GPUMask({ - id: `${this.id}-compose`, - inputs: this.predicates.map(predicate => predicate.mask), - output: finalMask - }).addToGraph(graph); +/** Adds the complete visibility workflow using one explicit dispatch limit. @internal */ +export function addGPUVisibilityWorkflowToGraphWithDispatchLimit( + workflow: GPUVisibilityWorkflow, + graph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + const template = workflow.predicates[0].mask; + for (const view of [ + ...workflow.predicates.flatMap(predicate => getVisibilityChunks(predicate.mask)), + ...getVisibilityChunks(workflow.output), + ...(workflow.outputMask ? getVisibilityChunks(workflow.outputMask) : []), + ...(workflow.sourceIds ? getVisibilityChunks(workflow.sourceIds) : []), + workflow.count + ]) { + if (view.buffer.graph !== graph) { + throw new Error(`${workflow.id} views must belong to the target graph`); } + } - const sourceIds = - this.sourceIds ?? createTransientVisibilityInput(graph, `${this.id}-source-ids`, template); - if (!this.sourceIds) { - addIdentityPasses(graph, `${this.id}-identity`, sourceIds, this.firstSourceIndex); - } + const finalMask = + workflow.outputMask ?? createTransientVisibilityInput(graph, `${workflow.id}-mask`, template); + if (finalMask !== template || workflow.predicates.length > 1) { + const mask = new GPUMask({ + id: `${workflow.id}-compose`, + inputs: workflow.predicates.map(predicate => predicate.mask), + output: finalMask + }); + addGPUMaskToGraphWithDispatchLimit(mask, graph, maxComputeWorkgroupsPerDimension); + } - new GPUCompaction({ - id: `${this.id}-compact`, - input: sourceIds, - flags: finalMask, - output: this.output, - count: this.count - }).addToGraph(graph); + const sourceIds = + workflow.sourceIds ?? + createTransientVisibilityInput(graph, `${workflow.id}-source-ids`, template); + if (!workflow.sourceIds) { + addIdentityPasses( + graph, + `${workflow.id}-identity`, + sourceIds, + workflow.firstSourceIndex, + maxComputeWorkgroupsPerDimension + ); } + + const compaction = new GPUCompaction({ + id: `${workflow.id}-compact`, + input: sourceIds, + flags: finalMask, + output: workflow.output, + count: workflow.count + }); + addGPUCompactionToGraphWithDispatchLimit(compaction, graph, maxComputeWorkgroupsPerDimension); } /** Creates graph-owned storage with the same atomic or vector topology as a visibility input. */ @@ -180,7 +211,8 @@ function addIdentityPasses( graph: GPUCommandGraph, id: string, output: GPUCompactionInput, - firstSourceIndex: number + firstSourceIndex: number, + maxComputeWorkgroupsPerDimension: number ): void { let chunkSourceOffset = firstSourceIndex; for (const [chunkIndex, chunk] of getVisibilityChunks(output).entries()) { @@ -188,7 +220,13 @@ function addIdentityPasses( addIdentityPass(graph, { id: output instanceof GraphVectorView ? `${id}-chunk-${chunkIndex}` : id, output: chunk, - firstSourceIndex: chunkSourceOffset + firstSourceIndex: chunkSourceOffset, + dispatchLayout: getBoundedDispatchLayout( + 'GPUVisibilityWorkflow', + chunk.length, + VISIBILITY_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ) }); } chunkSourceOffset += chunk.length; @@ -198,7 +236,12 @@ function addIdentityPasses( /** Writes consecutive uint32 source IDs into one packed view. */ function addIdentityPass( graph: GPUCommandGraph, - props: {id: string; output: GraphDataView<'uint32'>; firstSourceIndex: number} + props: { + id: string; + output: GraphDataView<'uint32'>; + firstSourceIndex: number; + dispatchLayout: GPUBoundedDispatchLayout; + } ): void { const source = /* wgsl */ ` const ELEMENT_COUNT: u32 = ${props.output.length}u; @@ -207,8 +250,11 @@ const FIRST_SOURCE_INDEX: u32 = ${props.firstSourceIndex}u; @group(0) @binding(0) var outputIds: array; @compute @workgroup_size(${VISIBILITY_WORKGROUP_SIZE}) -fn main(@builtin(global_invocation_id) globalId: vec3) { - let index = globalId.x; +fn main( + @builtin(local_invocation_index) localInvocationIndex: u32, + @builtin(workgroup_id) workgroupId: vec3 +) { + ${getBoundedInvocationIndexSource(props.dispatchLayout, VISIBILITY_WORKGROUP_SIZE)} if (index < ELEMENT_COUNT) { outputIds[OUTPUT_OFFSET + index] = FIRST_SOURCE_INDEX + index; } @@ -232,7 +278,9 @@ fn main(@builtin(global_invocation_id) globalId: vec3) { computation.setBindings(bindings); computation.dispatch( computePass, - Math.ceil(props.output.length / VISIBILITY_WORKGROUP_SIZE) + props.dispatchLayout.x, + props.dispatchLayout.y, + props.dispatchLayout.z ); }, destroy: () => computation.destroy() diff --git a/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.node.spec.ts b/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.node.spec.ts new file mode 100644 index 0000000000..269bada0ed --- /dev/null +++ b/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.node.spec.ts @@ -0,0 +1,139 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// Copyright (c) vis.gl contributors + +import {Buffer} from '@luma.gl/core'; +import {GPUCommandGraph, GPUCompaction, GPUVisibilityWorkflow} from '@luma.gl/experimental'; +import {NullDevice} from '@luma.gl/test-utils'; +import {describe, expect, test, vi} from 'vitest'; +import {addGPUCompactionToGraphWithDispatchLimit} from '../../src/gpu-primitives/gpu-compaction'; +import { + getBoundedDispatchLayout, + getBoundedInvocationIndexSource +} from '../../src/gpu-primitives/gpu-dispatch-utils'; +import {addGPUVisibilityWorkflowToGraphWithDispatchLimit} from '../../src/gpu-primitives/gpu-visibility-workflow'; + +const WORKGROUP_SIZE = 256; + +describe('bounded GPU visibility dispatch', () => { + test('preserves exact one-dimensional boundaries and expands safely into the third dimension', () => { + const maximum = 65_535; + const oneDimensionalCapacity = maximum * WORKGROUP_SIZE; + + for (const operation of ['GPUMask', 'GPUVisibilityWorkflow', 'GPUCompaction']) { + expect(getBoundedDispatchLayout(operation, 0, WORKGROUP_SIZE, maximum)).toEqual({ + x: 1, + y: 1, + z: 1 + }); + expect( + getBoundedDispatchLayout(operation, oneDimensionalCapacity, WORKGROUP_SIZE, maximum) + ).toEqual({x: maximum, y: 1, z: 1}); + expect( + getBoundedDispatchLayout(operation, oneDimensionalCapacity + 1, WORKGROUP_SIZE, maximum) + ).toEqual({x: maximum, y: 2, z: 1}); + expect( + getBoundedDispatchLayout(operation, 4 * WORKGROUP_SIZE + 1, WORKGROUP_SIZE, 2) + ).toEqual({x: 2, y: 2, z: 2}); + expect(() => + getBoundedDispatchLayout(operation, 8 * WORKGROUP_SIZE + 1, WORKGROUP_SIZE, 2) + ).toThrow(/exceeding the 3D dispatch limit/i); + } + + const source = getBoundedInvocationIndexSource({x: 2, y: 2, z: 2}, WORKGROUP_SIZE); + expect(source).toContain('workgroupId.z * 2u + workgroupId.y'); + expect(source).toContain('* 2u + workgroupId.x'); + expect(source.indexOf('workgroupIndex >= 16777216u')).toBeLessThan( + source.indexOf('workgroupIndex * 256u + localInvocationIndex') + ); + }); + + test('propagates a synthetic dispatch limit through every visibility workflow stage', () => { + const fixture = createVisibilityGraphFixture(4 * WORKGROUP_SIZE + 1); + const addComputePass = vi.spyOn(fixture.graph, 'addComputePass'); + + try { + addGPUVisibilityWorkflowToGraphWithDispatchLimit(fixture.workflow, fixture.graph, 2); + + const passIds = addComputePass.mock.calls.map(([pass]) => pass.id); + expect(passIds).toContain('node-visibility-compose'); + expect(passIds).toContain('node-visibility-identity'); + expect(passIds).toContain('node-visibility-compact-scan-level-0-scan'); + expect(passIds).toContain('node-visibility-compact-scan-level-0-add-offsets'); + expect(passIds).toContain('node-visibility-compact-scatter'); + } finally { + addComputePass.mockRestore(); + fixture.device.destroy(); + } + }); + + test('rejects source ranges beyond the complete bounded three-dimensional capacity', () => { + const fixture = createVisibilityGraphFixture(8 * WORKGROUP_SIZE + 1); + + try { + expect(() => + addGPUVisibilityWorkflowToGraphWithDispatchLimit(fixture.workflow, fixture.graph, 2) + ).toThrow(/GPUMask.*exceeding the 3D dispatch limit/i); + } finally { + fixture.device.destroy(); + } + }); + + test('retains the existing single-workgroup clear-count path for empty compaction', () => { + const device = new NullDevice({id: 'empty-compaction-node-device'}); + Object.defineProperty(device, 'type', {value: 'webgpu'}); + const graph = new GPUCommandGraph(device, {id: 'empty-compaction-node-graph'}); + const input = createTransientView(graph, 'input', 0); + const flags = createTransientView(graph, 'flags', 0); + const output = createTransientView(graph, 'output', 0); + const count = createTransientView(graph, 'count', 1); + const addComputePass = vi.spyOn(graph, 'addComputePass'); + + try { + const compaction = new GPUCompaction({id: 'empty-compaction', input, flags, output, count}); + addGPUCompactionToGraphWithDispatchLimit(compaction, graph, 0); + expect(addComputePass.mock.calls.map(([pass]) => pass.id)).toEqual([ + 'empty-compaction-clear-count' + ]); + } finally { + addComputePass.mockRestore(); + device.destroy(); + } + }); +}); + +function createVisibilityGraphFixture(rowCount: number): { + device: NullDevice; + graph: GPUCommandGraph; + workflow: GPUVisibilityWorkflow; +} { + const device = new NullDevice({id: 'bounded-visibility-node-device'}); + Object.defineProperty(device, 'type', {value: 'webgpu'}); + const graph = new GPUCommandGraph(device, {id: 'bounded-visibility-node-graph'}); + const firstMask = createTransientView(graph, 'first-mask', rowCount); + const secondMask = createTransientView(graph, 'second-mask', rowCount); + const outputMask = createTransientView(graph, 'output-mask', rowCount); + const output = createTransientView(graph, 'output', rowCount); + const count = createTransientView(graph, 'count', 1); + const workflow = new GPUVisibilityWorkflow({ + id: 'node-visibility', + predicates: [ + {kind: 'bounds', mask: firstMask}, + {kind: 'selection', mask: secondMask} + ], + outputMask, + output, + count, + firstSourceIndex: 40 + }); + return {device, graph, workflow}; +} + +function createTransientView(graph: GPUCommandGraph, id: string, length: number) { + const buffer = graph.createTransientBuffer({ + id, + byteLength: Math.max(length, 1) * Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE + }); + return graph.createDataView(buffer, {format: 'uint32', length}); +} diff --git a/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.spec.ts b/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.spec.ts index 710f82369d..589e4a2006 100644 --- a/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.spec.ts +++ b/modules/experimental/test/gpu-primitives/gpu-visibility-workflow.spec.ts @@ -4,6 +4,7 @@ import test from 'test/utils/vitest-tape'; import {Buffer, type Device} from '@luma.gl/core'; +import {Computation} from '@luma.gl/engine'; import { DrawCommandBuffer, GPUCommandGraph, @@ -12,6 +13,8 @@ import { } from '@luma.gl/experimental'; import {GPUData, GPUVector} from '@luma.gl/tables'; import {getWebGPUTestDevice} from '@luma.gl/test-utils'; +import {vi} from 'vitest'; +import {addGPUVisibilityWorkflowToGraphWithDispatchLimit} from '../../src/gpu-primitives/gpu-visibility-workflow'; test('GPUVisibilityWorkflow composes predicates and publishes indirect-ready results', async t => { const device = await getWebGPUTestDevice(); @@ -138,6 +141,120 @@ test('GPUVisibilityWorkflow composes predicates and publishes indirect-ready res t.end(); }); +test('GPUVisibilityWorkflow scales mask, identity, scan, and scatter through bounded 3D dispatches', async t => { + const device = await getWebGPUTestDevice(); + if (!device) { + t.comment('WebGPU is not available'); + t.end(); + return; + } + + const rowCount = 4 * 256 + 1; + const firstSourceIndex = 40; + const firstPredicate = Uint32Array.from({length: rowCount}, (_, index) => + index % 3 === 0 ? 2 : 0 + ); + const secondPredicate = Uint32Array.from({length: rowCount}, (_, index) => + index % 5 === 0 ? 0 : 4 + ); + const expectedMask = Array.from(firstPredicate, (value, index) => + value !== 0 && secondPredicate[index] !== 0 ? 1 : 0 + ); + const expectedSourceIds = expectedMask.flatMap((selected, index) => + selected ? [firstSourceIndex + index] : [] + ); + const firstPredicateBuffer = device.createBuffer({ + id: 'bounded-visibility-first-predicate', + data: firstPredicate, + usage: Buffer.STORAGE | Buffer.COPY_DST + }); + const secondPredicateBuffer = device.createBuffer({ + id: 'bounded-visibility-second-predicate', + data: secondPredicate, + usage: Buffer.STORAGE | Buffer.COPY_DST + }); + const outputMaskBuffer = device.createBuffer({ + id: 'bounded-visibility-output-mask', + byteLength: rowCount * Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE | Buffer.COPY_SRC + }); + const outputBuffer = device.createBuffer({ + id: 'bounded-visibility-output', + byteLength: rowCount * Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE | Buffer.COPY_SRC + }); + const countBuffer = device.createBuffer({ + id: 'bounded-visibility-count', + byteLength: Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE | Buffer.COPY_SRC + }); + const graph = new GPUCommandGraph(device, {id: 'bounded-visibility-graph'}); + const importView = (id: string, buffer: Buffer, length: number) => + graph.importGPUData(id, new GPUData({buffer, format: 'uint32', length, ownsBuffer: false})); + const workflow = new GPUVisibilityWorkflow({ + id: 'bounded-visibility', + predicates: [ + {kind: 'bounds', mask: importView('first-predicate', firstPredicateBuffer, rowCount)}, + {kind: 'selection', mask: importView('second-predicate', secondPredicateBuffer, rowCount)} + ], + outputMask: importView('output-mask', outputMaskBuffer, rowCount), + output: importView('output', outputBuffer, rowCount), + count: importView('count', countBuffer, 1), + firstSourceIndex + }); + addGPUVisibilityWorkflowToGraphWithDispatchLimit(workflow, graph, 2); + const compiled = graph.compile(); + const dispatchSpy = vi.spyOn(Computation.prototype, 'dispatch'); + + try { + await encodeAndSubmit(device, compiled, 'bounded-visibility-encoding'); + + t.deepEqual( + await readUint32(outputMaskBuffer, rowCount), + expectedMask, + 'multidimensional mask composition visits each source row exactly once' + ); + t.deepEqual( + await readUint32(outputBuffer, expectedSourceIds.length), + expectedSourceIds, + 'bounded identity generation and scatter preserve stable source-row order' + ); + t.deepEqual( + await readUint32(countBuffer, 1), + [expectedSourceIds.length], + 'padded multidimensional workgroups never inflate the selected count' + ); + + const dispatches = dispatchSpy.mock.instances.map((computation, index) => ({ + id: (computation as Computation).id, + dimensions: dispatchSpy.mock.calls[index].slice(1) + })); + for (const passId of [ + 'bounded-visibility-compose', + 'bounded-visibility-identity', + 'bounded-visibility-compact-scan-level-0-scan', + 'bounded-visibility-compact-scan-level-0-add-offsets', + 'bounded-visibility-compact-scatter' + ]) { + t.deepEqual( + dispatches.find(dispatch => dispatch.id === passId)?.dimensions, + [2, 2, 2], + `${passId} inherits the same bounded three-dimensional dispatch limit` + ); + } + } finally { + dispatchSpy.mockRestore(); + compiled.destroy(); + firstPredicateBuffer.destroy(); + secondPredicateBuffer.destroy(); + outputMaskBuffer.destroy(); + outputBuffer.destroy(); + countBuffer.destroy(); + } + + t.end(); +}); + test('GPUVisibilityWorkflow preserves chunk topology while generating global IDs', async t => { const device = await getWebGPUTestDevice(); if (!device) {