From 8e6ec66ead546fee2b4acd32f1a87c52e7ebbfad Mon Sep 17 00:00:00 2001 From: Ib Green Date: Tue, 4 Aug 2026 21:23:21 -0400 Subject: [PATCH] feat(experimental): compute GPU weakly connected components --- modules/experimental/src/lugraph/index.ts | 2 + ...lu-graph-connected-components-internals.ts | 374 ++++++++++ .../lugraph/lu-graph-connected-components.ts | 165 +++++ ...lu-graph-connected-components.node.spec.ts | 457 ++++++++++++ .../lu-graph-connected-components.spec.ts | 653 ++++++++++++++++++ 5 files changed, 1651 insertions(+) create mode 100644 modules/experimental/src/lugraph/lu-graph-connected-components-internals.ts create mode 100644 modules/experimental/src/lugraph/lu-graph-connected-components.ts create mode 100644 modules/experimental/test/lugraph/lu-graph-connected-components.node.spec.ts create mode 100644 modules/experimental/test/lugraph/lu-graph-connected-components.spec.ts diff --git a/modules/experimental/src/lugraph/index.ts b/modules/experimental/src/lugraph/index.ts index b831e2f14f..2d1c2a71c6 100644 --- a/modules/experimental/src/lugraph/index.ts +++ b/modules/experimental/src/lugraph/index.ts @@ -14,3 +14,5 @@ export type { LuGraphBreadthFirstSearchDirection, LuGraphBreadthFirstSearchProps } from './lu-graph-breadth-first-search'; +export {LuGraphConnectedComponents} from './lu-graph-connected-components'; +export type {LuGraphConnectedComponentsProps} from './lu-graph-connected-components'; diff --git a/modules/experimental/src/lugraph/lu-graph-connected-components-internals.ts b/modules/experimental/src/lugraph/lu-graph-connected-components-internals.ts new file mode 100644 index 0000000000..2b85f7f329 --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-connected-components-internals.ts @@ -0,0 +1,374 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: Copyright (c) vis.gl contributors +// SPDX-FileComment: Independently implemented for WebGPU; inspired by NVIDIA RAPIDS cuGraph. + +import {type Binding} from '@luma.gl/core'; +import {Computation} from '@luma.gl/engine'; +import type { + GPUCommandGraph, + GraphBufferUse, + GraphDataView +} from '../gpu-primitives/gpu-command-graph'; +import { + type GPUBoundedDispatchLayout, + getBoundedDispatchLayout, + getBoundedInvocationIndexSource +} from '../gpu-primitives/gpu-dispatch-utils'; +import {getViewBinding, getViewElementOffset} from '../gpu-primitives/graph-data-view-utils'; +import type {LuGraphConnectedComponents} from './lu-graph-connected-components'; + +const CONNECTED_COMPONENTS_WORKGROUP_SIZE = 256; +const INVALID_COMPONENT = 0xffffffff; + +type ImportedConnectedComponents = { + id: string; + vertexCount: number; + offsets: GraphDataView<'uint32'>; + neighbors: GraphDataView<'uint32'>; + overflow: GraphDataView<'uint32'>; + output: GraphDataView<'uint32'>; + converged?: GraphDataView<'uint32'>; + maxComputeWorkgroupsPerDimension: number; +}; + +type ConnectedComponentsBinding = { + view: GraphDataView<'uint32'>; + usage: GraphBufferUse['usage']; + atomic?: boolean; +}; + +type ConnectedComponentsPassProps = { + id: string; + source: string; + bindings: Record; + dispatchLayout: GPUBoundedDispatchLayout; +}; + +/** Adds bounded GPU weak-component hooking using an explicit dispatch limit. @internal */ +export function addLuGraphConnectedComponentsToGraphWithDispatchLimit( + components: LuGraphConnectedComponents, + commandGraph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + if (components.topology.graph.vertexCount === 0 && !components.converged) { + return; + } + + const state: ImportedConnectedComponents = { + id: components.id, + vertexCount: components.topology.graph.vertexCount, + offsets: commandGraph.importGPUVector( + `${components.id}-offsets`, + components.topology.forward.offsets + ).data[0], + neighbors: commandGraph.importGPUVector( + `${components.id}-neighbors`, + components.topology.forward.neighbors + ).data[0], + overflow: commandGraph.importGPUVector( + `${components.id}-overflow`, + components.topology.forward.overflow + ).data[0], + output: commandGraph.importGPUVector(`${components.id}-output`, components.output).data[0], + ...(components.converged + ? { + converged: commandGraph.importGPUVector( + `${components.id}-converged`, + components.converged + ).data[0] + } + : {}), + maxComputeWorkgroupsPerDimension + }; + + addInitializationPass(commandGraph, state); + if (state.vertexCount === 0) { + return; + } + + for (let iteration = 0; iteration < components.iterations; iteration++) { + if (state.converged) { + addConvergenceResetPass(commandGraph, {state, iteration}); + } + addHookingPass(commandGraph, {state, iteration}); + addPointerJumpPass(commandGraph, {state, iteration}); + } +} + +/** Initializes identity labels, fails closed on overflow, and marks empty graphs converged. */ +function addInitializationPass( + commandGraph: GPUCommandGraph, + state: ImportedConnectedComponents +): void { + const bindings: Record = { + output: {view: state.output, usage: 'storage-write', atomic: true}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.converged + ? {converged: {view: state.converged, usage: 'storage-write', atomic: true}} + : {}) + }; + const convergenceOffset = state.converged + ? `const CONVERGED_OFFSET: u32 = ${getViewElementOffset(state.converged)}u;` + : ''; + const convergenceInitialization = state.converged + ? `if (index == 0u) { + let emptyAndValid = VERTEX_COUNT == 0u && !hasOverflow; + atomicStore(&converged[CONVERGED_OFFSET], select(0u, 1u, emptyAndValid)); + }` + : ''; + const dispatchLayout = getLuGraphConnectedComponentsDispatchLayout( + Math.max(state.vertexCount, 1), + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${convergenceOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${CONNECTED_COMPONENTS_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, CONNECTED_COMPONENTS_WORKGROUP_SIZE)} + let hasOverflow = overflow[OVERFLOW_OFFSET] != 0u; + if (index < VERTEX_COUNT) { + let component = select(index, ${INVALID_COMPONENT}u, hasOverflow); + atomicStore(&output[OUTPUT_OFFSET + index], component); + } + ${convergenceInitialization} +}`; + + addConnectedComponentsPass(commandGraph, { + id: `${state.id}-initialize`, + source, + bindings, + dispatchLayout + }); +} + +/** Resets the caller-owned convergence scalar in a dedicated globally synchronized pass. */ +function addConvergenceResetPass( + commandGraph: GPUCommandGraph, + props: {state: ImportedConnectedComponents; iteration: number} +): void { + const {state} = props; + const converged = state.converged!; + const bindings: Record = { + overflow: {view: state.overflow, usage: 'storage-read'}, + converged: {view: converged, usage: 'storage-write', atomic: true} + }; + const dispatchLayout = getLuGraphConnectedComponentsDispatchLayout( + 1, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +const CONVERGED_OFFSET: u32 = ${getViewElementOffset(converged)}u; +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${CONNECTED_COMPONENTS_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, CONNECTED_COMPONENTS_WORKGROUP_SIZE)} + if (index != 0u) { return; } + atomicStore(&converged[CONVERGED_OFFSET], select(1u, 0u, overflow[OVERFLOW_OFFSET] != 0u)); +}`; + + addConnectedComponentsPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-reset`, + source, + bindings, + dispatchLayout + }); +} + +/** Hooks both endpoints of every forward edge into one order-independent weak component. */ +function addHookingPass( + commandGraph: GPUCommandGraph, + props: {state: ImportedConnectedComponents; iteration: number} +): void { + const {state} = props; + const bindings: Record = { + offsets: {view: state.offsets, usage: 'storage-read'}, + neighbors: {view: state.neighbors, usage: 'storage-read'}, + output: {view: state.output, usage: 'storage-read-write', atomic: true}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.converged + ? {converged: {view: state.converged, usage: 'storage-read-write', atomic: true}} + : {}) + }; + const convergenceOffset = state.converged + ? `const CONVERGED_OFFSET: u32 = ${getViewElementOffset(state.converged)}u;` + : ''; + const markChanged = state.converged + ? `if (previous > lower) { atomicStore(&converged[CONVERGED_OFFSET], 0u); }` + : ''; + const dispatchLayout = getLuGraphConnectedComponentsDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const CAPACITY: u32 = ${state.neighbors.length}u; +const OFFSETS_OFFSET: u32 = ${getViewElementOffset(state.offsets)}u; +const NEIGHBORS_OFFSET: u32 = ${getViewElementOffset(state.neighbors)}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${convergenceOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${CONNECTED_COMPONENTS_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, CONNECTED_COMPONENTS_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT || overflow[OVERFLOW_OFFSET] != 0u) { return; } + let first = min(offsets[OFFSETS_OFFSET + index], CAPACITY); + let last = min(offsets[OFFSETS_OFFSET + index + 1u], CAPACITY); + for (var slot = first; slot < last; slot++) { + let neighbor = neighbors[NEIGHBORS_OFFSET + slot]; + if (neighbor >= VERTEX_COUNT) { continue; } + let sourceComponent = atomicLoad(&output[OUTPUT_OFFSET + index]); + let neighborComponent = atomicLoad(&output[OUTPUT_OFFSET + neighbor]); + let lower = min(sourceComponent, neighborComponent); + let higher = max(sourceComponent, neighborComponent); + if (higher == lower) { continue; } + let previous = atomicMin(&output[OUTPUT_OFFSET + higher], lower); + ${markChanged} + } +}`; + + addConnectedComponentsPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-hook`, + source, + bindings, + dispatchLayout + }); +} + +/** Shortcuts the monotone parent forest in a pass synchronized after all edge hooks. */ +function addPointerJumpPass( + commandGraph: GPUCommandGraph, + props: {state: ImportedConnectedComponents; iteration: number} +): void { + const {state} = props; + const bindings: Record = { + output: {view: state.output, usage: 'storage-read-write', atomic: true}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.converged + ? {converged: {view: state.converged, usage: 'storage-read-write', atomic: true}} + : {}) + }; + const convergenceOffset = state.converged + ? `const CONVERGED_OFFSET: u32 = ${getViewElementOffset(state.converged)}u;` + : ''; + const markChanged = state.converged + ? `if (previous > parentComponent) { + atomicStore(&converged[CONVERGED_OFFSET], 0u); + }` + : ''; + const dispatchLayout = getLuGraphConnectedComponentsDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${convergenceOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${CONNECTED_COMPONENTS_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, CONNECTED_COMPONENTS_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT || overflow[OVERFLOW_OFFSET] != 0u) { return; } + let component = atomicLoad(&output[OUTPUT_OFFSET + index]); + let parentComponent = atomicLoad(&output[OUTPUT_OFFSET + component]); + if (parentComponent >= component) { return; } + let previous = atomicMin(&output[OUTPUT_OFFSET + index], parentComponent); + ${markChanged} +}`; + + addConnectedComponentsPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-jump`, + source, + bindings, + dispatchLayout + }); +} + +/** Declares every storage binding in its exact generated shader-layout order. */ +function getBindingDeclarations(bindings: Record): string { + return Object.entries(bindings) + .map(([name, binding], location) => { + const access = binding.usage === 'storage-read' ? 'read' : 'read_write'; + const element = binding.atomic ? 'atomic' : 'u32'; + return `@group(0) @binding(${location}) var ${name}: array<${element}>;`; + }) + .join('\n'); +} + +/** Compiles one bounded storage-buffer pass without hidden GPU allocation or submission. */ +function addConnectedComponentsPass( + commandGraph: GPUCommandGraph, + props: ConnectedComponentsPassProps +): void { + commandGraph.addComputePass({ + id: props.id, + resources: Object.values(props.bindings).map(({view, usage}) => ({buffer: view, usage})), + compile: ({device}) => { + const computation = new Computation(device, { + id: props.id, + source: props.source, + shaderLayout: { + bindings: Object.keys(props.bindings).map((name, location) => ({ + name, + type: 'storage' as const, + group: 0, + location + })) + } + }); + + return { + encode: ({computePass, getBuffer}) => { + const bindings: Record = {}; + for (const [name, binding] of Object.entries(props.bindings)) { + bindings[name] = getViewBinding(binding.view, getBuffer); + } + computation.setBindings(bindings); + computation.dispatch( + computePass, + props.dispatchLayout.x, + props.dispatchLayout.y, + props.dispatchLayout.z + ); + }, + destroy: () => computation.destroy() + }; + } + }); +} + +/** Plans a bounded three-dimensional weak-component vertex or status dispatch. @internal */ +export function getLuGraphConnectedComponentsDispatchLayout( + elementCount: number, + maxComputeWorkgroupsPerDimension: number +): GPUBoundedDispatchLayout { + return getBoundedDispatchLayout( + 'LuGraphConnectedComponents', + elementCount, + CONNECTED_COMPONENTS_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ); +} diff --git a/modules/experimental/src/lugraph/lu-graph-connected-components.ts b/modules/experimental/src/lugraph/lu-graph-connected-components.ts new file mode 100644 index 0000000000..6e7bd353db --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-connected-components.ts @@ -0,0 +1,165 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: Copyright (c) vis.gl contributors +// SPDX-FileComment: Independently implemented for WebGPU; inspired by NVIDIA RAPIDS cuGraph. + +import type {Buffer} from '@luma.gl/core'; +import {DynamicBuffer} from '@luma.gl/engine'; +import type {GPUData, GPUVector} from '@luma.gl/tables'; +import type {GPUCommandGraph} from '../gpu-primitives/gpu-command-graph'; +import {addLuGraphConnectedComponentsToGraphWithDispatchLimit} from './lu-graph-connected-components-internals'; +import type {LuGraphAdjacency, LuGraphTopology} from './lu-graph-topology'; + +const DEFAULT_COMPONENT_ITERATIONS = 32; +const MAXIMUM_COMPONENT_ITERATIONS = 1024; +const SCALAR_BYTE_LENGTH = 4; + +/** Caller-owned GPU graph topology, component labels, and optional convergence status. */ +export type LuGraphConnectedComponentsProps = { + /** Prefix for generated command-graph node and imported-resource identifiers. */ + id?: string; + /** Existing forward graph adjacency; reverse adjacency is not required. */ + topology: LuGraphTopology; + /** One caller-owned packed unsigned component label for every graph vertex. */ + output: GPUVector<'uint32'>; + /** Bounded number of compiled relaxation and pointer-jumping iterations. Defaults to 32. */ + iterations?: number; + /** Optional caller-owned scalar publishing whether the final iteration reached a fixed point. */ + converged?: GPUVector<'uint32'>; +}; + +/** + * Publishes weak graph components directly from existing GPU-resident forward adjacency. + * + * Directed edges connect both endpoints, so reverse adjacency is unnecessary. Once converged, each + * component receives its lowest stable vertex identifier, and isolated vertices label themselves. + * The optional convergence scalar is one only when the last compiled iteration reaches a fixed + * point. Forward-adjacency overflow instead publishes `0xffffffff` labels and zero convergence. + */ +export class LuGraphConnectedComponents { + /** Prefix for generated command-graph node and imported-resource identifiers. */ + readonly id: string; + /** Existing caller-owned GPU graph topology. */ + readonly topology: LuGraphTopology; + /** Caller-owned vertex-aligned component labels. */ + readonly output: GPUVector<'uint32'>; + /** Number of compiled, explicitly synchronized component-relaxation iterations. */ + readonly iterations: number; + /** Optional caller-owned GPU-resident convergence status. */ + readonly converged?: GPUVector<'uint32'>; + + /** Validates caller-owned graph metadata without allocating, submitting, or reading GPU work. */ + constructor(props: LuGraphConnectedComponentsProps) { + this.id = props.id ?? 'lu-graph-connected-components'; + this.topology = props.topology; + this.output = props.output; + this.iterations = props.iterations ?? DEFAULT_COMPONENT_ITERATIONS; + this.converged = props.converged; + + if ( + !Number.isSafeInteger(this.iterations) || + this.iterations < 1 || + this.iterations > MAXIMUM_COMPONENT_ITERATIONS + ) { + throw new Error(`${this.id} iterations must be a safe integer between one and 1024`); + } + validateComponentVector(this.output, this.topology.graph.vertexCount, `${this.id} output`); + if (this.converged) { + validateComponentVector(this.converged, 1, `${this.id} converged`); + } + validateDistinctComponentOutputs(this); + } + + /** Declares bounded weak-component passes without submitting commands or reading results. */ + addToGraph(commandGraph: GPUCommandGraph): void { + addLuGraphConnectedComponentsToGraphWithDispatchLimit( + this, + commandGraph, + commandGraph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} + +/** Requires one packed, aligned unsigned output chunk with its exact logical row count. */ +function validateComponentVector(vector: GPUVector<'uint32'>, length: number, name: string): void { + if ( + vector.data.length !== 1 || + vector.format !== 'uint32' || + vector.stride !== 1 || + vector.byteStride !== SCALAR_BYTE_LENGTH || + vector.rowByteLength !== SCALAR_BYTE_LENGTH || + vector.valueLength !== vector.length || + vector.bufferLayout + ) { + throw new Error(`${name} must contain exactly one packed uint32 chunk`); + } + if (vector.length !== length) { + throw new Error(`${name} must contain exactly ${length} uint32 rows`); + } + + const chunk = vector.data[0]; + if ( + chunk.format !== 'uint32' || + chunk.length !== length || + chunk.stride !== 1 || + chunk.byteStride !== SCALAR_BYTE_LENGTH || + chunk.rowByteLength !== SCALAR_BYTE_LENGTH || + chunk.valueLength !== chunk.length || + !Number.isSafeInteger(chunk.byteOffset) || + chunk.byteOffset < 0 || + chunk.byteOffset % SCALAR_BYTE_LENGTH !== 0 + ) { + throw new Error(`${name} must contain one packed, uint32-aligned chunk`); + } +} + +/** Keeps component labels and optional convergence status disjoint from all graph allocations. */ +function validateDistinctComponentOutputs(components: LuGraphConnectedComponents): void { + const topology = components.topology; + const inputVectors = [ + topology.graph.sourceVertices, + topology.graph.targetVertices, + ...(topology.graph.edgeWeights ? [topology.graph.edgeWeights] : []), + ...(topology.graph.edgeIds ? [topology.graph.edgeIds] : []), + ...getAdjacencyVectors(topology.forward), + ...(topology.reverse ? getAdjacencyVectors(topology.reverse) : []), + topology.invalidEdgeCount + ]; + const allocations = new Set(); + for (const vector of inputVectors) { + for (const chunk of vector.data) { + allocations.add(getPhysicalBuffer(chunk)); + } + } + + const outputs = [ + {name: 'output', vector: components.output}, + ...(components.converged ? [{name: 'converged', vector: components.converged}] : []) + ]; + for (const {name, vector} of outputs) { + const buffer = getPhysicalBuffer(vector.data[0]); + if (allocations.has(buffer)) { + throw new Error(`${components.id} ${name} must use a distinct physical buffer allocation`); + } + allocations.add(buffer); + } +} + +/** Enumerates existing caller-owned adjacency and status columns without changing their chunks. */ +function getAdjacencyVectors( + adjacency: LuGraphAdjacency +): (GPUVector<'uint32'> | GPUVector<'float32'>)[] { + return [ + adjacency.offsets, + adjacency.neighbors, + adjacency.edgeIds, + ...(adjacency.edgeWeights ? [adjacency.edgeWeights] : []), + adjacency.count, + adjacency.overflow + ]; +} + +/** Resolves a stable engine wrapper to its current concrete physical GPU allocation. */ +function getPhysicalBuffer(chunk: GPUData<'uint32'> | GPUData<'float32'>): Buffer { + return chunk.buffer instanceof DynamicBuffer ? chunk.buffer.buffer : chunk.buffer; +} diff --git a/modules/experimental/test/lugraph/lu-graph-connected-components.node.spec.ts b/modules/experimental/test/lugraph/lu-graph-connected-components.node.spec.ts new file mode 100644 index 0000000000..ac1d05eee2 --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-connected-components.node.spec.ts @@ -0,0 +1,457 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: Copyright (c) vis.gl contributors + +import {Buffer} from '@luma.gl/core'; +import {DynamicBuffer} from '@luma.gl/engine'; +import * as experimentalModule from '@luma.gl/experimental'; +import { + LuGraph, + LuGraphConnectedComponents, + LuGraphTopology, + type LuGraphAdjacency, + type LuGraphConnectedComponentsProps +} from '@luma.gl/experimental/lugraph'; +import {GPUData, GPUVector} from '@luma.gl/tables'; +import {NullDevice} from '@luma.gl/test-utils'; +import {afterEach, describe, expect, test, vi} from 'vitest'; + +type ScalarFormat = 'uint32' | 'float32'; +type ScalarValues = Uint32Array | Float32Array; + +type ComponentsFixture = { + device: NullDevice; + buffers: Buffer[]; + dynamicBuffers: DynamicBuffer[]; + vectors: GPUVector[]; +}; + +type VectorOptions = { + buffer?: Buffer | DynamicBuffer; + byteOffset?: number; + byteStride?: number; + rowByteLength?: number; + stride?: number; +}; + +const componentsFixtures: ComponentsFixture[] = []; + +afterEach(() => { + vi.restoreAllMocks(); + for (const fixture of componentsFixtures.splice(0)) { + for (const vector of fixture.vectors) vector.destroy(); + for (const dynamicBuffer of fixture.dynamicBuffers) dynamicBuffer.destroy(); + for (const buffer of fixture.buffers) buffer.destroy(); + fixture.device.destroy(); + } +}); + +describe('LuGraphConnectedComponents public contract and ownership', () => { + test('keeps weak connected components isolated in the optional luGraph entry point', () => { + expect(typeof LuGraphConnectedComponents).toBe('function'); + expect('LuGraphConnectedComponents' in experimentalModule).toBe(false); + }); + + test('preserves graph topology and caller outputs without allocations, submission, or readback', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {reverse: true, weighted: true, status: true}); + const createBufferSpy = vi.spyOn(fixture.device, 'createBuffer'); + const createCommandEncoderSpy = vi.spyOn(fixture.device, 'createCommandEncoder'); + const submitSpy = vi.spyOn(fixture.device, 'submit'); + const readbackSpies = fixture.buffers.map(buffer => vi.spyOn(buffer, 'readAsync')); + + const components = new LuGraphConnectedComponents({...props, id: 'borrowed-components'}); + + expect(components.id).toBe('borrowed-components'); + expect(components.topology).toBe(props.topology); + expect(components.output).toBe(props.output); + expect(components.converged).toBe(props.converged); + expect(components.iterations).toBe(32); + expect(components.topology.graph.sourceVertices.data.map(chunk => chunk.length)).toEqual([ + 2, 0, 3 + ]); + expect(createBufferSpy).not.toHaveBeenCalled(); + expect(createCommandEncoderSpy).not.toHaveBeenCalled(); + expect(submitSpy).not.toHaveBeenCalled(); + for (const readbackSpy of readbackSpies) expect(readbackSpy).not.toHaveBeenCalled(); + expect(Reflect.has(components, 'destroy')).toBe(false); + + for (const vector of fixture.vectors) vector.destroy(); + expect(fixture.buffers.every(buffer => !buffer.destroyed)).toBe(true); + }); + + test('accepts a directed graph without reverse adjacency and without optional convergence output', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const components = new LuGraphConnectedComponents(props); + + expect(components.topology.graph.directed).toBe(true); + expect(components.topology.reverse).toBeUndefined(); + expect(components.converged).toBeUndefined(); + }); + + test('accepts empty graph outputs and optional uint32 convergence status', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {vertexCount: 0, status: true}); + const components = new LuGraphConnectedComponents(props); + + expect(components.output.length).toBe(0); + expect(components.output.data).toHaveLength(1); + expect(components.output.data[0].buffer.byteLength).toBeGreaterThanOrEqual(4); + expect(components.converged?.length).toBe(1); + }); +}); + +describe('LuGraphConnectedComponents iteration and vector validation', () => { + test.each([1, 32, 1024])('accepts a positive bounded iteration count: %i', iterations => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + + expect(new LuGraphConnectedComponents({...props, iterations}).iterations).toBe(iterations); + }); + + test.each([ + 0, + -1, + 1.5, + Number.NaN, + Number.POSITIVE_INFINITY, + 1025 + ])('rejects an invalid or excessive iteration count: %s', iterations => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + + expect(() => new LuGraphConnectedComponents({...props, iterations})).toThrow( + /iterations|positive|1024/ + ); + }); + + test.each([5, 7])('requires exactly one output row per graph vertex: %i', length => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const output = createVector(fixture, `component-length-${length}`, 'uint32', [ + new Uint32Array(length) + ]); + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /output|vertexCount|length/ + ); + }); + + test('requires uint32 component IDs instead of float32', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const output = createVector(fixture, 'float-components', 'float32', [ + new Float32Array(props.topology.graph.vertexCount) + ]) as unknown as GPUVector<'uint32'>; + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /output|uint32|packed/ + ); + }); + + test.each([0, 2])('requires exactly one physical component-ID chunk: %i', chunkCount => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const chunks = chunkCount === 0 ? [] : [new Uint32Array(3), new Uint32Array(3)]; + const output = createVector(fixture, 'partitioned-components', 'uint32', chunks); + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /output|one|single|chunk/ + ); + }); + + test.each([ + ['misaligned byte offset', {byteOffset: 2}], + ['padded byte stride', {byteStride: 8}], + ['oversized row payload', {rowByteLength: 8}], + ['multi-component scalar stride', {stride: 2}] + ] as [string, VectorOptions][])('rejects unpacked component output: %s', (_name, options) => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const output = createVector( + fixture, + 'unpacked-components', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + options + ); + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /output|packed|aligned|uint32/ + ); + }); + + test.each([0, 2])('requires exactly one convergence status row: %i', length => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {status: true}); + const converged = createVector(fixture, `convergence-length-${length}`, 'uint32', [ + new Uint32Array(length) + ]); + + expect(() => new LuGraphConnectedComponents({...props, converged})).toThrow( + /converged|one|row|scalar/ + ); + }); + + test('requires packed uint32 convergence status with exactly one chunk', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {status: true}); + const wrongFormat = createVector(fixture, 'float-convergence', 'float32', [ + new Float32Array(1) + ]) as unknown as GPUVector<'uint32'>; + const partitioned = createVector(fixture, 'partitioned-convergence', 'uint32', [ + new Uint32Array(1), + new Uint32Array(0) + ]); + + expect(() => new LuGraphConnectedComponents({...props, converged: wrongFormat})).toThrow( + /converged|uint32|packed/ + ); + expect(() => new LuGraphConnectedComponents({...props, converged: partitioned})).toThrow( + /converged|one|single|chunk/ + ); + }); + + test('accepts uint32-aligned component and status ranges at non-256-byte offsets', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const output = createVector( + fixture, + 'offset-components', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + {byteOffset: 4} + ); + const converged = createVector(fixture, 'offset-convergence', 'uint32', [new Uint32Array(1)], { + byteOffset: 4 + }); + + const components = new LuGraphConnectedComponents({...props, output, converged}); + expect(components.output.data[0].byteOffset).toBe(4); + expect(components.converged?.data[0].byteOffset).toBe(4); + }); + + test.each([ + 'sourceVertices', + 'targetVertices', + 'edgeWeights', + 'edgeIds', + 'forward.offsets', + 'forward.neighbors', + 'forward.edgeIds', + 'forward.edgeWeights', + 'forward.count', + 'forward.overflow', + 'reverse.offsets', + 'reverse.neighbors', + 'reverse.edgeIds', + 'reverse.edgeWeights', + 'reverse.count', + 'reverse.overflow', + 'invalidEdgeCount' + ])('rejects writable labels backed by existing physical allocation: %s', vectorName => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {vertexCount: 1, reverse: true, weighted: true}); + const vector = getTopologyVector(props.topology, vectorName); + const output = createVector(fixture, 'aliased-components', 'uint32', [new Uint32Array(1)], { + buffer: vector.data[0].buffer + }); + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /output|distinct|physical|allocation/ + ); + }); + + test('rejects convergence status aliasing component labels or topology allocations', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture, {vertexCount: 1, status: true}); + const outputAlias = createVector( + fixture, + 'aliased-output-status', + 'uint32', + [new Uint32Array(1)], + { + buffer: props.output.data[0].buffer + } + ); + const topologyAlias = createVector( + fixture, + 'aliased-topology-status', + 'uint32', + [new Uint32Array(1)], + {buffer: props.topology.invalidEdgeCount.data[0].buffer} + ); + + expect(() => new LuGraphConnectedComponents({...props, converged: outputAlias})).toThrow( + /converged|distinct|physical|allocation/ + ); + expect(() => new LuGraphConnectedComponents({...props, converged: topologyAlias})).toThrow( + /converged|distinct|physical|allocation/ + ); + }); + + test('unwraps borrowed DynamicBuffer wrappers when checking physical component aliases', () => { + const fixture = createComponentsFixture(); + const props = createComponentsProps(fixture); + const concreteBuffer = props.topology.forward.offsets.data[0].buffer as Buffer; + const dynamicBuffer = new DynamicBuffer(fixture.device, { + id: 'borrowed-offset-wrapper', + buffer: concreteBuffer, + ownsBuffer: false + }); + fixture.dynamicBuffers.push(dynamicBuffer); + const output = createVector( + fixture, + 'dynamic-aliased-components', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + {buffer: dynamicBuffer} + ); + + expect(() => new LuGraphConnectedComponents({...props, output})).toThrow( + /distinct|physical|allocation/ + ); + expect(concreteBuffer.destroyed).toBe(false); + }); +}); + +function createComponentsFixture(): ComponentsFixture { + const fixture = {device: new NullDevice({}), buffers: [], dynamicBuffers: [], vectors: []}; + componentsFixtures.push(fixture); + return fixture; +} + +function createComponentsProps( + fixture: ComponentsFixture, + options: { + vertexCount?: number; + directed?: boolean; + reverse?: boolean; + weighted?: boolean; + status?: boolean; + } = {} +): LuGraphConnectedComponentsProps { + const vertexCount = options.vertexCount ?? 6; + const sourceVertices = createVector(fixture, 'sourceVertices', 'uint32', [ + Uint32Array.from([0, 2]), + new Uint32Array(0), + Uint32Array.from([2, 3, 4]) + ]); + const targetVertices = createVector(fixture, 'targetVertices', 'uint32', [ + Uint32Array.from([1, 4]), + new Uint32Array(0), + Uint32Array.from([3, 5, 1]) + ]); + const edgeWeights = options.weighted + ? createVector(fixture, 'sourceWeights', 'float32', [ + Float32Array.from([0.5, 2]), + new Float32Array(0), + Float32Array.from([1, 4, 8]) + ]) + : undefined; + const edgeIds = options.weighted + ? createVector(fixture, 'sourceEdgeIds', 'uint32', [ + Uint32Array.from([10, 20]), + new Uint32Array(0), + Uint32Array.from([30, 40, 50]) + ]) + : undefined; + const graph = new LuGraph({ + vertexCount, + sourceVertices, + targetVertices, + edgeWeights, + edgeIds, + directed: options.directed + }); + const forward = createAdjacency(fixture, 'forward', vertexCount, 5, options.weighted); + const reverse = options.reverse + ? createAdjacency(fixture, 'reverse', vertexCount, 5, options.weighted) + : undefined; + const invalidEdgeCount = createVector(fixture, 'invalidEdgeCount', 'uint32', [ + new Uint32Array(1) + ]); + const topology = new LuGraphTopology({graph, forward, reverse, invalidEdgeCount}); + const output = createVector(fixture, 'componentIds', 'uint32', [new Uint32Array(vertexCount)]); + const converged = options.status + ? createVector(fixture, 'converged', 'uint32', [new Uint32Array(1)]) + : undefined; + return {topology, output, converged}; +} + +function createAdjacency( + fixture: ComponentsFixture, + name: string, + vertexCount: number, + capacity: number, + weighted = false +): LuGraphAdjacency { + return { + offsets: createVector(fixture, `${name}-offsets`, 'uint32', [new Uint32Array(vertexCount + 1)]), + neighbors: createVector(fixture, `${name}-neighbors`, 'uint32', [new Uint32Array(capacity)]), + edgeIds: createVector(fixture, `${name}-edgeIds`, 'uint32', [new Uint32Array(capacity)]), + edgeWeights: weighted + ? createVector(fixture, `${name}-weights`, 'float32', [new Float32Array(capacity)]) + : undefined, + count: createVector(fixture, `${name}-count`, 'uint32', [new Uint32Array(1)]), + overflow: createVector(fixture, `${name}-overflow`, 'uint32', [new Uint32Array(1)]) + }; +} + +function getTopologyVector(topology: LuGraphTopology, name: string): GPUVector { + if (name === 'invalidEdgeCount') return topology.invalidEdgeCount; + if (name === 'sourceVertices') return topology.graph.sourceVertices; + if (name === 'targetVertices') return topology.graph.targetVertices; + if (name === 'edgeWeights') return topology.graph.edgeWeights!; + if (name === 'edgeIds') return topology.graph.edgeIds!; + + const [direction, vectorName] = name.split('.'); + const adjacency = direction === 'forward' ? topology.forward : topology.reverse!; + return adjacency[vectorName as keyof LuGraphAdjacency]!; +} + +function createVector( + fixture: ComponentsFixture, + name: string, + format: Format, + chunks: readonly ScalarValues[], + options: VectorOptions = {} +): GPUVector { + const byteOffset = options.byteOffset ?? 0; + const byteStride = options.byteStride ?? Uint32Array.BYTES_PER_ELEMENT; + const rowByteLength = options.rowByteLength ?? Uint32Array.BYTES_PER_ELEMENT; + const stride = options.stride ?? 1; + const data = chunks.map((values, chunkIndex) => { + const buffer = + options.buffer ?? + fixture.device.createBuffer({ + id: `${name}-chunk-${chunkIndex}-${fixture.buffers.length}`, + byteLength: byteOffset + Math.max(Math.max(values.length, 1) * byteStride, rowByteLength), + usage: Buffer.STORAGE | Buffer.COPY_DST | Buffer.COPY_SRC + }); + if (!options.buffer) fixture.buffers.push(buffer as Buffer); + return new GPUData({ + buffer, + format, + length: values.length, + byteOffset, + byteStride, + rowByteLength, + stride, + ownsBuffer: false + }); + }); + const vector = new GPUVector({ + type: 'data', + name, + format, + data, + byteStride, + rowByteLength, + stride, + ownsData: false + }); + fixture.vectors.push(vector); + return vector; +} diff --git a/modules/experimental/test/lugraph/lu-graph-connected-components.spec.ts b/modules/experimental/test/lugraph/lu-graph-connected-components.spec.ts new file mode 100644 index 0000000000..f1d446fc8b --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-connected-components.spec.ts @@ -0,0 +1,653 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: Copyright (c) vis.gl contributors + +import {Buffer, type Device} from '@luma.gl/core'; +import {GPUCommandGraph} from '@luma.gl/experimental'; +import { + LuGraph, + LuGraphConnectedComponents, + LuGraphTopology, + type LuGraphAdjacency +} from '@luma.gl/experimental/lugraph'; +import {GPUData, GPUVector} from '@luma.gl/tables'; +import {getWebGPUTestDevice} from '@luma.gl/test-utils'; +import test, {type Test} from 'test/utils/vitest-tape'; +import {vi} from 'vitest'; +import { + addLuGraphConnectedComponentsToGraphWithDispatchLimit, + getLuGraphConnectedComponentsDispatchLayout +} from '../../src/lugraph/lu-graph-connected-components-internals'; + +const INVALID_COMPONENT = 0xffffffff; + +type ScalarFormat = 'uint32' | 'float32'; + +type ComponentsScenario = { + name: string; + vertexCount: number; + sourceChunks: number[][]; + targetChunks: number[][]; + weightChunks?: number[][]; + directed?: boolean; + reverse?: boolean; + capacity?: number; + reverseCapacity?: number; + iterations?: number; + status?: boolean; + incomplete?: boolean; + allowPartialLabels?: boolean; + maximumWorkgroups?: number; + byteOffset?: number; + assertNoScratch?: boolean; +}; + +type ExpectedComponents = { + labels: number[]; + invalidEdgeCount: number; + forwardCount: number; + reverseCount: number; + forwardOverflow: boolean; + reverseOverflow: boolean; +}; + +type ComponentsExecutionFixture = { + device: Device; + buffers: Buffer[]; + vectors: GPUVector[]; + graph: LuGraph; + topology: LuGraphTopology; + components: LuGraphConnectedComponents; + commandGraph: GPUCommandGraph; + compiled?: ReturnType; +}; + +const componentsScenarios: ComponentsScenario[] = [ + { + name: 'empty graphs publish a converged status and preserve zero-length label ownership', + vertexCount: 0, + sourceChunks: [], + targetChunks: [], + capacity: 0 + }, + { + name: 'empty graphs without convergence status do not require component work', + vertexCount: 0, + sourceChunks: [], + targetChunks: [], + capacity: 0, + status: false + }, + { + name: 'isolated vertices retain their own stable component identifiers', + vertexCount: 7, + sourceChunks: [[], []], + targetChunks: [[], []], + capacity: 0, + iterations: 1 + }, + { + name: 'directed edges form weak components without requiring reverse CSR adjacency', + vertexCount: 6, + sourceChunks: [[3, 2], [], [1]], + targetChunks: [[2, 1], [], [0]], + iterations: 6, + assertNoScratch: true + }, + { + name: 'disconnected directed components select their lowest stable vertex identifier', + vertexCount: 8, + sourceChunks: [[4, 3], [], [6]], + targetChunks: [[1, 2], [], [5]], + iterations: 4 + }, + { + name: 'cycles, duplicate edges, diamonds, and self-loops converge deterministically', + vertexCount: 8, + sourceChunks: [[4, 3, 3], [], [1, 2, 2, 6]], + targetChunks: [[3, 2, 2], [], [2, 4, 2, 6]], + iterations: 8 + }, + { + name: 'weighted undirected chunks preserve minimum IDs while ignoring edge weights', + vertexCount: 7, + sourceChunks: [[0, 3], [], [3, 5]], + targetChunks: [[1, 2], [], [4, 5]], + weightChunks: [[0.5, 2], [], [4, 8]], + directed: false, + iterations: 5 + }, + { + name: 'invalid endpoints are excluded without connecting unrelated isolated vertices', + vertexCount: 6, + sourceChunks: [[0, 9], [], [2, 4, 5]], + targetChunks: [[1, 2], [], [8, 5, 5]], + iterations: 5 + }, + { + name: 'optional convergence output can be omitted while labels remain deterministic', + vertexCount: 5, + sourceChunks: [[4, 2]], + targetChunks: [[2, 1]], + iterations: 5, + status: false + }, + { + name: 'reverse adjacency overflow is irrelevant to forward-only weak connectivity', + vertexCount: 5, + sourceChunks: [[0, 1, 3]], + targetChunks: [[1, 2, 4]], + reverse: true, + reverseCapacity: 0, + iterations: 5 + }, + { + name: 'zero forward capacity fails closed with invalid labels and unconverged status', + vertexCount: 4, + sourceChunks: [[0, 1]], + targetChunks: [[1, 2]], + capacity: 0, + iterations: 4 + }, + { + name: 'partial forward adjacency also fails closed without exposing partial components', + vertexCount: 5, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + capacity: 2, + iterations: 4 + }, + { + name: 'a final changed iteration conservatively reports incomplete even if labels are canonical', + vertexCount: 2, + sourceChunks: [[0]], + targetChunks: [[1]], + iterations: 1, + incomplete: true + }, + { + name: 'a bounded long-chain iteration exposes only valid monotone partial component labels', + vertexCount: 12, + sourceChunks: [[11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1]], + targetChunks: [[10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0]], + iterations: 1, + incomplete: true, + allowPartialLabels: true + }, + { + name: 'a no-change final iteration proves convergence after component hooking', + vertexCount: 2, + sourceChunks: [[0]], + targetChunks: [[1]], + iterations: 2 + }, + { + name: 'non-256-aligned CSR offsets, component labels, and convergence status stay correct', + vertexCount: 5, + sourceChunks: [[3, 1]], + targetChunks: [[2, 0]], + iterations: 4, + byteOffset: 4 + }, + { + name: 'bounded three-dimensional hooking reaches the final of 1025 vertices', + vertexCount: 1025, + sourceChunks: [[1024]], + targetChunks: [[0]], + iterations: 2, + maximumWorkgroups: 2 + } +]; + +test('LuGraphConnectedComponents plans bounded three-dimensional vertex dispatch', tapeTest => { + tapeTest.deepEqual(getLuGraphConnectedComponentsDispatchLayout(0, 2), {x: 1, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphConnectedComponentsDispatchLayout(512, 2), {x: 2, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphConnectedComponentsDispatchLayout(513, 2), {x: 2, y: 2, z: 1}); + tapeTest.deepEqual(getLuGraphConnectedComponentsDispatchLayout(1025, 2), {x: 2, y: 2, z: 2}); + tapeTest.throws(() => getLuGraphConnectedComponentsDispatchLayout(2049, 2), /3D dispatch limit/); + tapeTest.end(); +}); + +for (const scenario of componentsScenarios) { + test(`LuGraphConnectedComponents GPU labeling: ${scenario.name}`, async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const expected = calculateExpectedComponents(scenario); + const fixture = createExecutionFixture(device, scenario, expected); + try { + compileComponents(fixture, scenario.maximumWorkgroups); + executeComponents(fixture); + await assertComponents(tapeTest, fixture, scenario, expected); + tapeTest.deepEqual( + fixture.graph.sourceVertices.data.map(chunk => chunk.length), + scenario.sourceChunks.map(chunk => chunk.length), + 'weak-component evaluation preserves every original source chunk' + ); + if (scenario.assertNoScratch) { + tapeTest.equal( + fixture.compiled?.stats.logicalTransientBufferCount, + 3, + 'weak-component hooking allocates no scratch beyond CSR topology construction' + ); + } + } finally { + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); + }); +} + +test('LuGraphConnectedComponents rebuilds weak labels after source updates without hidden execution', async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const original: ComponentsScenario = { + name: 'repeat weak components', + vertexCount: 6, + sourceChunks: [[0, 1], [], [3, 4]], + targetChunks: [[1, 2], [], [4, 5]], + iterations: 5 + }; + const fixture = createExecutionFixture(device, original, calculateExpectedComponents(original)); + const submitSpy = vi.spyOn(device, 'submit'); + const sourceReadbackSpies = [ + ...fixture.graph.sourceVertices.data, + ...fixture.graph.targetVertices.data + ].map(chunk => vi.spyOn(chunk.buffer, 'readAsync')); + + try { + compileComponents(fixture); + tapeTest.equal( + submitSpy.mock.calls.length, + 0, + 'construction and compilation never submit work' + ); + tapeTest.ok( + sourceReadbackSpies.every(spy => spy.mock.calls.length === 0), + 'weak connectivity never reads graph source buffers back' + ); + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + + executeComponents(fixture); + await assertComponents(tapeTest, fixture, original, calculateExpectedComponents(original)); + + const sourceBuffer = fixture.graph.sourceVertices.data[0].buffer as Buffer; + sourceBuffer.write(Uint32Array.from([9, 1])); + const updated = {...original, sourceChunks: [[9, 1], [], [3, 4]]}; + executeComponents(fixture); + await assertComponents(tapeTest, fixture, updated, calculateExpectedComponents(updated)); + tapeTest.equal( + fixture.graph.sourceVertices.data[0].buffer, + sourceBuffer, + 'source updates retain exact caller-owned chunk identity' + ); + } finally { + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); +}); + +/** Computes weak connected components with minimum stable vertex IDs using CPU union-find. */ +function calculateExpectedComponents(scenario: ComponentsScenario): ExpectedComponents { + const parents = Array.from({length: scenario.vertexCount}, (_, vertexIndex) => vertexIndex); + let invalidEdgeCount = 0; + let validEdgeCount = 0; + let forwardCount = 0; + + for (const [chunkIndex, sources] of scenario.sourceChunks.entries()) { + for (const [rowIndex, source] of sources.entries()) { + const target = scenario.targetChunks[chunkIndex][rowIndex]; + if (source >= scenario.vertexCount || target >= scenario.vertexCount) { + invalidEdgeCount++; + continue; + } + validEdgeCount++; + forwardCount += scenario.directed === false && source !== target ? 2 : 1; + mergeComponents(parents, source, target); + } + } + + const labels = parents.map((_parent, vertexIndex) => findComponentRoot(parents, vertexIndex)); + const forwardOverflow = forwardCount > (scenario.capacity ?? forwardCount); + const reverseCount = scenario.reverse ? validEdgeCount : 0; + const reverseOverflow = reverseCount > (scenario.reverseCapacity ?? reverseCount); + + return { + labels: forwardOverflow ? new Array(scenario.vertexCount).fill(INVALID_COMPONENT) : labels, + invalidEdgeCount, + forwardCount, + reverseCount, + forwardOverflow, + reverseOverflow + }; +} + +function findComponentRoot(parents: number[], vertex: number): number { + let root = vertex; + while (parents[root] !== root) root = parents[root]; + return root; +} + +function mergeComponents(parents: number[], source: number, target: number): void { + const sourceRoot = findComponentRoot(parents, source); + const targetRoot = findComponentRoot(parents, target); + if (sourceRoot !== targetRoot) { + parents[Math.max(sourceRoot, targetRoot)] = Math.min(sourceRoot, targetRoot); + } +} + +function createExecutionFixture( + device: Device, + scenario: ComponentsScenario, + expected: ExpectedComponents +): ComponentsExecutionFixture { + const buffers: Buffer[] = []; + const vectors: GPUVector[] = []; + const sourceVertices = createInputVector( + device, + buffers, + vectors, + 'source-vertices', + 'uint32', + scenario.sourceChunks + ); + const targetVertices = createInputVector( + device, + buffers, + vectors, + 'target-vertices', + 'uint32', + scenario.targetChunks + ); + const edgeWeights = scenario.weightChunks + ? createInputVector( + device, + buffers, + vectors, + 'source-weights', + 'float32', + scenario.weightChunks + ) + : undefined; + const graph = new LuGraph({ + vertexCount: scenario.vertexCount, + sourceVertices, + targetVertices, + edgeWeights, + directed: scenario.directed + }); + const forward = createOutputAdjacency( + device, + buffers, + vectors, + 'forward', + scenario.vertexCount, + scenario.capacity ?? expected.forwardCount, + Boolean(edgeWeights), + scenario.byteOffset + ); + const reverse = scenario.reverse + ? createOutputAdjacency( + device, + buffers, + vectors, + 'reverse', + scenario.vertexCount, + scenario.reverseCapacity ?? expected.reverseCount, + Boolean(edgeWeights), + scenario.byteOffset + ) + : undefined; + const invalidEdgeCount = createOutputVector( + device, + buffers, + vectors, + 'invalid-edges', + 'uint32', + 1 + ); + const topology = new LuGraphTopology({graph, forward, reverse, invalidEdgeCount}); + const output = createOutputVector( + device, + buffers, + vectors, + 'component-identifiers', + 'uint32', + scenario.vertexCount, + scenario.byteOffset + ); + const converged = + scenario.status === false + ? undefined + : createOutputVector( + device, + buffers, + vectors, + 'components-converged', + 'uint32', + 1, + scenario.byteOffset + ); + const components = new LuGraphConnectedComponents({ + topology, + output, + iterations: scenario.iterations, + converged + }); + + return { + device, + buffers, + vectors, + graph, + topology, + components, + commandGraph: new GPUCommandGraph(device) + }; +} + +function createInputVector( + device: Device, + buffers: Buffer[], + vectors: GPUVector[], + name: string, + format: Format, + chunks: readonly number[][] +): GPUVector { + const data = chunks.map((chunk, chunkIndex) => { + const values = format === 'float32' ? Float32Array.from(chunk) : Uint32Array.from(chunk); + const buffer = device.createBuffer({ + id: `${name}-chunk-${chunkIndex}`, + data: values.length > 0 ? values : new Uint32Array(1), + usage: Buffer.STORAGE | Buffer.COPY_DST + }); + buffers.push(buffer); + return new GPUData({buffer, format, length: values.length, ownsBuffer: false}); + }); + const vector = new GPUVector({type: 'data', name, format, data, ownsData: false}); + vectors.push(vector); + return vector; +} + +function createOutputAdjacency( + device: Device, + buffers: Buffer[], + vectors: GPUVector[], + name: string, + vertexCount: number, + capacity: number, + weighted: boolean, + byteOffset = 0 +): LuGraphAdjacency { + return { + offsets: createOutputVector( + device, + buffers, + vectors, + `${name}-offsets`, + 'uint32', + vertexCount + 1, + byteOffset + ), + neighbors: createOutputVector( + device, + buffers, + vectors, + `${name}-neighbors`, + 'uint32', + capacity + ), + edgeIds: createOutputVector(device, buffers, vectors, `${name}-edge-ids`, 'uint32', capacity), + edgeWeights: weighted + ? createOutputVector(device, buffers, vectors, `${name}-weights`, 'float32', capacity) + : undefined, + count: createOutputVector(device, buffers, vectors, `${name}-count`, 'uint32', 1), + overflow: createOutputVector(device, buffers, vectors, `${name}-overflow`, 'uint32', 1) + }; +} + +function createOutputVector( + device: Device, + buffers: Buffer[], + vectors: GPUVector[], + name: string, + format: Format, + length: number, + byteOffset = 0 +): GPUVector { + const buffer = device.createBuffer({ + id: name, + byteLength: byteOffset + Math.max(length, 1) * Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE | Buffer.COPY_SRC + }); + buffers.push(buffer); + const vector = new GPUVector({ + type: 'buffer', + name, + format, + buffer, + length, + byteOffset, + ownsBuffer: false + }); + vectors.push(vector); + return vector; +} + +function compileComponents(fixture: ComponentsExecutionFixture, maximumWorkgroups?: number): void { + fixture.topology.addToGraph(fixture.commandGraph); + if (maximumWorkgroups === undefined) { + fixture.components.addToGraph(fixture.commandGraph); + } else { + addLuGraphConnectedComponentsToGraphWithDispatchLimit( + fixture.components, + fixture.commandGraph, + maximumWorkgroups + ); + } + fixture.compiled = fixture.commandGraph.compile(); +} + +function executeComponents(fixture: ComponentsExecutionFixture): void { + const commandEncoder = fixture.device.createCommandEncoder({id: 'lu-graph-components-test'}); + fixture.compiled!.encode(commandEncoder, {parameters: undefined}); + fixture.device.submit(commandEncoder.finish()); +} + +async function assertComponents( + tapeTest: Test, + fixture: ComponentsExecutionFixture, + scenario: ComponentsScenario, + expected: ExpectedComponents +): Promise { + const [labels, convergence, invalidEdgeCount, forwardOverflow, reverseOverflow] = + await Promise.all([ + readUint32Vector(fixture.components.output), + fixture.components.converged + ? readUint32Vector(fixture.components.converged) + : Promise.resolve(undefined), + readUint32Vector(fixture.topology.invalidEdgeCount), + readUint32Vector(fixture.topology.forward.overflow), + fixture.topology.reverse + ? readUint32Vector(fixture.topology.reverse.overflow) + : Promise.resolve(undefined) + ]); + + if (scenario.allowPartialLabels && !expected.forwardOverflow) { + tapeTest.ok( + labels.every( + (label, vertexIndex) => + label <= vertexIndex && expected.labels[label] === expected.labels[vertexIndex] + ), + 'bounded iterations preserve monotone labels inside each true weak component' + ); + } else { + tapeTest.deepEqual(labels, expected.labels, 'component IDs equal each weak component minimum'); + } + + if (convergence) { + const expectedConvergence = Number(!expected.forwardOverflow && !scenario.incomplete); + tapeTest.equal( + convergence[0], + expectedConvergence, + 'convergence is proven only when the final iteration makes no changes' + ); + } + tapeTest.equal( + invalidEdgeCount[0], + expected.invalidEdgeCount, + 'invalid source edges stay excluded' + ); + tapeTest.equal( + forwardOverflow[0], + Number(expected.forwardOverflow), + 'forward adjacency overflow remains explicit' + ); + if (reverseOverflow) { + tapeTest.equal( + reverseOverflow[0], + Number(expected.reverseOverflow), + 'unused reverse overflow does not invalidate weak components' + ); + } + if (expected.forwardOverflow) { + tapeTest.ok( + labels.every(label => label === INVALID_COMPONENT), + 'truncated forward adjacency publishes no misleading partial labels' + ); + } +} + +async function readUint32Vector(vector: GPUVector<'uint32'>): Promise { + if (vector.length === 0) return []; + const data = vector.data[0]; + const bytes = await (data.buffer as Buffer).readAsync( + data.byteOffset, + vector.length * Uint32Array.BYTES_PER_ELEMENT + ); + return Array.from(new Uint32Array(bytes.buffer, bytes.byteOffset, vector.length)); +} + +function destroyExecutionFixture(tapeTest: Test, fixture: ComponentsExecutionFixture): void { + fixture.compiled?.destroy(); + for (const vector of fixture.vectors) vector.destroy(); + tapeTest.ok( + fixture.buffers.every(buffer => !buffer.destroyed), + 'compiled component graphs and borrowed vectors never destroy caller-owned buffers' + ); + for (const buffer of fixture.buffers) buffer.destroy(); +}