diff --git a/modules/experimental/src/lugraph/index.ts b/modules/experimental/src/lugraph/index.ts index 455fd5f7f5..b831e2f14f 100644 --- a/modules/experimental/src/lugraph/index.ts +++ b/modules/experimental/src/lugraph/index.ts @@ -9,3 +9,8 @@ export {LuGraphTopology} from './lu-graph-topology'; export type {LuGraphAdjacency, LuGraphTopologyProps} from './lu-graph-topology'; export {LuGraphDegree} from './lu-graph-degree'; export type {LuGraphDegreeDirection, LuGraphDegreeProps} from './lu-graph-degree'; +export {LuGraphBreadthFirstSearch} from './lu-graph-breadth-first-search'; +export type { + LuGraphBreadthFirstSearchDirection, + LuGraphBreadthFirstSearchProps +} from './lu-graph-breadth-first-search'; diff --git a/modules/experimental/src/lugraph/lu-graph-breadth-first-search-internals.ts b/modules/experimental/src/lugraph/lu-graph-breadth-first-search-internals.ts new file mode 100644 index 0000000000..9c262785a9 --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-breadth-first-search-internals.ts @@ -0,0 +1,417 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// Copyright (c) vis.gl contributors + +import {type Binding} from '@luma.gl/core'; +import {Computation} from '@luma.gl/engine'; +import type { + GPUCommandGraph, + GraphBufferUse, + GraphDataView, + GraphVectorView +} 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 {LuGraphBreadthFirstSearch} from './lu-graph-breadth-first-search'; +import type {LuGraphAdjacency} from './lu-graph-topology'; + +const BREADTH_FIRST_SEARCH_WORKGROUP_SIZE = 256; +const UNREACHABLE_VERTEX = 0xffffffff; + +type ImportedSearchAdjacency = { + offsets: GraphDataView<'uint32'>; + neighbors: GraphDataView<'uint32'>; + overflow: GraphDataView<'uint32'>; +}; + +type ImportedBreadthFirstSearch = { + id: string; + vertexCount: number; + seeds: GraphVectorView<'uint32'>; + distances: GraphDataView<'uint32'>; + predecessors: GraphDataView<'uint32'>; + mask?: GraphDataView<'uint32'>; + seedCount?: GraphDataView<'uint32'>; + activeDepth?: GraphDataView<'uint32'>; + primaryAdjacency: ImportedSearchAdjacency; + secondaryAdjacency?: ImportedSearchAdjacency; + maxComputeWorkgroupsPerDimension: number; +}; + +type BreadthFirstSearchBinding = { + view: GraphDataView<'uint32'>; + usage: GraphBufferUse['usage']; + atomic?: boolean; +}; + +type BreadthFirstSearchPassProps = { + id: string; + source: string; + bindings: Record; + dispatchLayout: GPUBoundedDispatchLayout; +}; + +/** Adds deterministic GPU shortest-hop traversal with an explicit dispatch limit. @internal */ +export function addLuGraphBreadthFirstSearchToGraphWithDispatchLimit( + search: LuGraphBreadthFirstSearch, + commandGraph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + if (search.topology.graph.vertexCount === 0) { + return; + } + + const useIncoming = search.topology.graph.directed && search.direction === 'incoming'; + const primaryAdjacency = importSearchAdjacency( + commandGraph, + `${search.id}-${useIncoming ? 'incoming' : 'outgoing'}`, + useIncoming ? search.topology.reverse! : search.topology.forward + ); + const state: ImportedBreadthFirstSearch = { + id: search.id, + vertexCount: search.topology.graph.vertexCount, + seeds: commandGraph.importGPUVector(`${search.id}-seeds`, search.seeds), + distances: commandGraph.importGPUVector(`${search.id}-distances`, search.distances).data[0], + predecessors: commandGraph.importGPUVector(`${search.id}-predecessors`, search.predecessors) + .data[0], + ...(search.mask + ? {mask: commandGraph.importGPUVector(`${search.id}-mask`, search.mask).data[0]} + : {}), + ...(search.seedCount + ? { + seedCount: commandGraph.importGPUVector(`${search.id}-seed-count`, search.seedCount) + .data[0] + } + : {}), + ...(search.activeDepth + ? { + activeDepth: commandGraph.importGPUVector(`${search.id}-active-depth`, search.activeDepth) + .data[0] + } + : {}), + primaryAdjacency, + ...(search.topology.graph.directed && search.direction === 'both' + ? { + secondaryAdjacency: importSearchAdjacency( + commandGraph, + `${search.id}-incoming`, + search.topology.reverse! + ) + } + : {}), + maxComputeWorkgroupsPerDimension + }; + + addInitializationPass(commandGraph, state); + + let seedBase = 0; + for (const [chunkIndex, seeds] of state.seeds.data.entries()) { + if (seeds.length > 0) { + addSeedPass(commandGraph, {state, seeds, seedBase, chunkIndex}); + } + seedBase += seeds.length; + } + + for (let depth = 0; depth < search.maxDepth; depth++) { + addExpansionPass(commandGraph, { + state, + adjacency: state.primaryAdjacency, + depth, + direction: useIncoming ? 'incoming' : 'outgoing' + }); + if (state.secondaryAdjacency) { + addExpansionPass(commandGraph, { + state, + adjacency: state.secondaryAdjacency, + depth, + direction: 'incoming' + }); + } + } +} + +/** Imports only the packed CSR allocations and overflow state consumed by one direction. */ +function importSearchAdjacency( + commandGraph: GPUCommandGraph, + id: string, + adjacency: LuGraphAdjacency +): ImportedSearchAdjacency { + return { + offsets: commandGraph.importGPUVector(`${id}-offsets`, adjacency.offsets).data[0], + neighbors: commandGraph.importGPUVector(`${id}-neighbors`, adjacency.neighbors).data[0], + overflow: commandGraph.importGPUVector(`${id}-overflow`, adjacency.overflow).data[0] + }; +} + +/** Resets published outputs before every encoding without allocating frontier scratch. */ +function addInitializationPass( + commandGraph: GPUCommandGraph, + state: ImportedBreadthFirstSearch +): void { + const bindings: Record = { + distances: {view: state.distances, usage: 'storage-write', atomic: true}, + predecessors: {view: state.predecessors, usage: 'storage-write', atomic: true}, + ...(state.mask ? {mask: {view: state.mask, usage: 'storage-write', atomic: true}} : {}) + }; + const maskOffset = state.mask + ? `const MASK_OFFSET: u32 = ${getViewElementOffset(state.mask)}u;` + : ''; + const clearMask = state.mask ? 'atomicStore(&mask[MASK_OFFSET + index], 0u);' : ''; + const dispatchLayout = getLuGraphBreadthFirstSearchDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const DISTANCES_OFFSET: u32 = ${getViewElementOffset(state.distances)}u; +const PREDECESSORS_OFFSET: u32 = ${getViewElementOffset(state.predecessors)}u; +${maskOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${BREADTH_FIRST_SEARCH_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, BREADTH_FIRST_SEARCH_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT) { return; } + atomicStore(&distances[DISTANCES_OFFSET + index], ${UNREACHABLE_VERTEX}u); + atomicStore(&predecessors[PREDECESSORS_OFFSET + index], ${UNREACHABLE_VERTEX}u); + ${clearMask} +}`; + + addBreadthFirstSearchPass(commandGraph, { + id: `${state.id}-initialize`, + source, + bindings, + dispatchLayout + }); +} + +/** Publishes valid roots from one original seed chunk unless selected adjacency overflowed. */ +function addSeedPass( + commandGraph: GPUCommandGraph, + props: { + state: ImportedBreadthFirstSearch; + seeds: GraphDataView<'uint32'>; + seedBase: number; + chunkIndex: number; + } +): void { + const {state} = props; + const bindings: Record = { + seeds: {view: props.seeds, usage: 'storage-read'}, + distances: {view: state.distances, usage: 'storage-read-write', atomic: true}, + overflow: {view: state.primaryAdjacency.overflow, usage: 'storage-read'}, + ...(state.secondaryAdjacency + ? {secondaryOverflow: {view: state.secondaryAdjacency.overflow, usage: 'storage-read'}} + : {}), + ...(state.mask ? {mask: {view: state.mask, usage: 'storage-read-write', atomic: true}} : {}), + ...(state.seedCount ? {activeSeedCount: {view: state.seedCount, usage: 'storage-read'}} : {}) + }; + const dynamicOffsets = [ + state.secondaryAdjacency + ? `const SECONDARY_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.secondaryAdjacency.overflow)}u;` + : '', + state.mask ? `const MASK_OFFSET: u32 = ${getViewElementOffset(state.mask)}u;` : '', + state.seedCount + ? `const ACTIVE_SEED_COUNT_OFFSET: u32 = ${getViewElementOffset(state.seedCount)}u;` + : '' + ].join('\n'); + const secondaryOverflowGuard = state.secondaryAdjacency + ? ' || secondaryOverflow[SECONDARY_OVERFLOW_OFFSET] != 0u' + : ''; + const seedCountGuard = state.seedCount + ? `if (${props.seedBase}u + index >= activeSeedCount[ACTIVE_SEED_COUNT_OFFSET]) { return; }` + : ''; + const publishMask = state.mask ? 'atomicStore(&mask[MASK_OFFSET + vertex], 1u);' : ''; + const dispatchLayout = getLuGraphBreadthFirstSearchDispatchLayout( + props.seeds.length, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const SEED_COUNT: u32 = ${props.seeds.length}u; +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const SEEDS_OFFSET: u32 = ${getViewElementOffset(props.seeds)}u; +const DISTANCES_OFFSET: u32 = ${getViewElementOffset(state.distances)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.primaryAdjacency.overflow)}u; +${dynamicOffsets} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${BREADTH_FIRST_SEARCH_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, BREADTH_FIRST_SEARCH_WORKGROUP_SIZE)} + if (index >= SEED_COUNT) { return; } + if (overflow[OVERFLOW_OFFSET] != 0u${secondaryOverflowGuard}) { return; } + ${seedCountGuard} + let vertex = seeds[SEEDS_OFFSET + index]; + if (vertex >= VERTEX_COUNT) { return; } + atomicStore(&distances[DISTANCES_OFFSET + vertex], 0u); + ${publishMask} +}`; + + addBreadthFirstSearchPass(commandGraph, { + id: `${state.id}-seed-${props.chunkIndex}`, + source, + bindings, + dispatchLayout + }); +} + +/** Computes the next shortest-hop layer and the deterministic lowest-ID predecessor. */ +function addExpansionPass( + commandGraph: GPUCommandGraph, + props: { + state: ImportedBreadthFirstSearch; + adjacency: ImportedSearchAdjacency; + depth: number; + direction: 'outgoing' | 'incoming'; + } +): void { + const {state, adjacency} = props; + const bindings: Record = { + offsets: {view: adjacency.offsets, usage: 'storage-read'}, + neighbors: {view: adjacency.neighbors, usage: 'storage-read'}, + distances: {view: state.distances, usage: 'storage-read-write', atomic: true}, + predecessors: {view: state.predecessors, usage: 'storage-read-write', atomic: true}, + overflow: {view: state.primaryAdjacency.overflow, usage: 'storage-read'}, + ...(state.secondaryAdjacency + ? {secondaryOverflow: {view: state.secondaryAdjacency.overflow, usage: 'storage-read'}} + : {}), + ...(state.mask ? {mask: {view: state.mask, usage: 'storage-read-write', atomic: true}} : {}), + ...(state.activeDepth ? {activeDepth: {view: state.activeDepth, usage: 'storage-read'}} : {}) + }; + const dynamicOffsets = [ + state.secondaryAdjacency + ? `const SECONDARY_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.secondaryAdjacency.overflow)}u;` + : '', + state.mask ? `const MASK_OFFSET: u32 = ${getViewElementOffset(state.mask)}u;` : '', + state.activeDepth + ? `const ACTIVE_DEPTH_OFFSET: u32 = ${getViewElementOffset(state.activeDepth)}u;` + : '' + ].join('\n'); + const secondaryOverflowGuard = state.secondaryAdjacency + ? ' || secondaryOverflow[SECONDARY_OVERFLOW_OFFSET] != 0u' + : ''; + const activeDepthGuard = state.activeDepth + ? `if (${props.depth}u >= activeDepth[ACTIVE_DEPTH_OFFSET]) { return; }` + : ''; + const publishMask = state.mask ? 'atomicStore(&mask[MASK_OFFSET + neighbor], 1u);' : ''; + const dispatchLayout = getLuGraphBreadthFirstSearchDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const CAPACITY: u32 = ${adjacency.neighbors.length}u; +const OFFSETS_OFFSET: u32 = ${getViewElementOffset(adjacency.offsets)}u; +const NEIGHBORS_OFFSET: u32 = ${getViewElementOffset(adjacency.neighbors)}u; +const DISTANCES_OFFSET: u32 = ${getViewElementOffset(state.distances)}u; +const PREDECESSORS_OFFSET: u32 = ${getViewElementOffset(state.predecessors)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.primaryAdjacency.overflow)}u; +${dynamicOffsets} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${BREADTH_FIRST_SEARCH_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, BREADTH_FIRST_SEARCH_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT) { return; } + if (overflow[OVERFLOW_OFFSET] != 0u${secondaryOverflowGuard}) { return; } + ${activeDepthGuard} + if (atomicLoad(&distances[DISTANCES_OFFSET + index]) != ${props.depth}u) { 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 previousDistance = atomicMin(&distances[DISTANCES_OFFSET + neighbor], ${props.depth + 1}u); + if (previousDistance >= ${props.depth + 1}u) { + atomicMin(&predecessors[PREDECESSORS_OFFSET + neighbor], index); + ${publishMask} + } + } +}`; + + addBreadthFirstSearchPass(commandGraph, { + id: `${state.id}-depth-${props.depth}-${props.direction}`, + source, + bindings, + dispatchLayout + }); +} + +/** Declares storage buffers in the same order as the generated shader binding layout. */ +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 GPU pass without allocating, submitting, or reading graph resources. */ +function addBreadthFirstSearchPass( + commandGraph: GPUCommandGraph, + props: BreadthFirstSearchPassProps +): 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 one bounded three-dimensional breadth-first seed or vertex dispatch. @internal */ +export function getLuGraphBreadthFirstSearchDispatchLayout( + elementCount: number, + maxComputeWorkgroupsPerDimension: number +): GPUBoundedDispatchLayout { + return getBoundedDispatchLayout( + 'LuGraphBreadthFirstSearch', + elementCount, + BREADTH_FIRST_SEARCH_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ); +} diff --git a/modules/experimental/src/lugraph/lu-graph-breadth-first-search.ts b/modules/experimental/src/lugraph/lu-graph-breadth-first-search.ts new file mode 100644 index 0000000000..0ef0ea378c --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-breadth-first-search.ts @@ -0,0 +1,255 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// Copyright (c) vis.gl contributors + +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 {addLuGraphBreadthFirstSearchToGraphWithDispatchLimit} from './lu-graph-breadth-first-search-internals'; +import type {LuGraphAdjacency, LuGraphTopology} from './lu-graph-topology'; + +const MAXIMUM_UINT32 = 0xffffffff; +const MAXIMUM_BREADTH_FIRST_DEPTH = 1024; +const SCALAR_BYTE_LENGTH = 4; + +/** Forward, reverse, or combined orientation for shortest unweighted graph paths. */ +export type LuGraphBreadthFirstSearchDirection = 'outgoing' | 'incoming' | 'both'; + +/** Caller-owned graph, seeds, GPU-resident controls, and shortest-path destinations. */ +export type LuGraphBreadthFirstSearchProps = { + /** Prefix for generated command-graph nodes and imported-resource identifiers. */ + id?: string; + /** Existing GPU-resident graph adjacency and its explicit overflow status. */ + topology: LuGraphTopology; + /** Stable source vertex identifiers, preserving every original seed chunk. */ + seeds: GPUVector<'uint32'>; + /** One caller-owned unsigned hop-distance row for every graph vertex. */ + distances: GPUVector<'uint32'>; + /** One caller-owned unsigned predecessor row for every graph vertex. */ + predecessors: GPUVector<'uint32'>; + /** Optional caller-owned vertex-aligned zero/one reachability mask. */ + mask?: GPUVector<'uint32'>; + /** Optional caller-owned scalar limiting the number of active seed rows. */ + seedCount?: GPUVector<'uint32'>; + /** Compiled maximum hop count, bounded by 1,024. Defaults to one. */ + maxDepth?: number; + /** Optional caller-owned scalar dynamically limiting the compiled hop count. */ + activeDepth?: GPUVector<'uint32'>; + /** Forward, reverse, or combined edge orientation. Defaults to outgoing. */ + direction?: LuGraphBreadthFirstSearchDirection; +}; + +/** + * Publishes deterministic shortest unweighted paths over existing GPU-resident adjacency. + * + * Unreachable vertices and seed predecessors contain `0xffffffff`. Equal-length predecessor ties + * select the lowest stable parent vertex identifier. Invalid or inactive seeds are ignored. Any + * overflow in the selected adjacency direction leaves every distance and predecessor unreachable + * and clears the optional mask, preventing partial topology from producing misleading paths. + */ +export class LuGraphBreadthFirstSearch { + /** Prefix for generated command-graph nodes and imported-resource identifiers. */ + readonly id: string; + /** Existing caller-owned graph topology. */ + readonly topology: LuGraphTopology; + /** Existing caller-owned, chunk-preserving seed identifiers. */ + readonly seeds: GPUVector<'uint32'>; + /** Caller-owned vertex-aligned hop distances. */ + readonly distances: GPUVector<'uint32'>; + /** Caller-owned vertex-aligned, deterministically selected predecessors. */ + readonly predecessors: GPUVector<'uint32'>; + /** Optional caller-owned vertex-aligned reachability mask. */ + readonly mask?: GPUVector<'uint32'>; + /** Optional GPU-resident active seed count. */ + readonly seedCount?: GPUVector<'uint32'>; + /** Number of compiled breadth-first expansion passes. */ + readonly maxDepth: number; + /** Optional GPU-resident active traversal depth. */ + readonly activeDepth?: GPUVector<'uint32'>; + /** Forward, reverse, or combined edge orientation. */ + readonly direction: LuGraphBreadthFirstSearchDirection; + + /** Validates caller-owned metadata without allocating, submitting, or reading GPU work. */ + constructor(props: LuGraphBreadthFirstSearchProps) { + this.id = props.id ?? 'lu-graph-breadth-first-search'; + this.topology = props.topology; + this.seeds = props.seeds; + this.distances = props.distances; + this.predecessors = props.predecessors; + this.mask = props.mask; + this.seedCount = props.seedCount; + this.maxDepth = props.maxDepth ?? 1; + this.activeDepth = props.activeDepth; + this.direction = props.direction ?? 'outgoing'; + + if (!['outgoing', 'incoming', 'both'].includes(this.direction)) { + throw new Error(`${this.id} direction must be outgoing, incoming, or both`); + } + if (this.direction !== 'outgoing' && this.topology.graph.directed && !this.topology.reverse) { + throw new Error( + `${this.id} incoming or bidirectional directed search requires reverse adjacency` + ); + } + if ( + !Number.isSafeInteger(this.maxDepth) || + this.maxDepth < 0 || + this.maxDepth > MAXIMUM_BREADTH_FIRST_DEPTH + ) { + throw new Error(`${this.id} maxDepth must be a safe integer between zero and 1024`); + } + + validateSeedVector(this.seeds, this.id); + validateSingleChunkVector( + this.distances, + this.topology.graph.vertexCount, + `${this.id} distances` + ); + validateSingleChunkVector( + this.predecessors, + this.topology.graph.vertexCount, + `${this.id} predecessors` + ); + if (this.mask) { + validateSingleChunkVector(this.mask, this.topology.graph.vertexCount, `${this.id} mask`); + } + if (this.seedCount) { + validateSingleChunkVector(this.seedCount, 1, `${this.id} seedCount`); + } + if (this.activeDepth) { + validateSingleChunkVector(this.activeDepth, 1, `${this.id} activeDepth`); + } + validateDistinctSearchOutputs(this); + } + + /** Adds bounded shortest-path passes without submitting commands or reading GPU results. */ + addToGraph(commandGraph: GPUCommandGraph): void { + addLuGraphBreadthFirstSearchToGraphWithDispatchLimit( + this, + commandGraph, + commandGraph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} + +/** Preserves ordered source seed chunks while validating their packed uint32 scalar layout. */ +function validateSeedVector(seeds: GPUVector<'uint32'>, id: string): void { + validatePackedVector(seeds, `${id} seeds`); + if ( + !Number.isSafeInteger(seeds.length) || + seeds.length < 0 || + seeds.length > MAXIMUM_UINT32 || + seeds.data.reduce((totalLength, chunk) => totalLength + chunk.length, 0) !== seeds.length + ) { + throw new Error(`${id} seed row count must fit in uint32`); + } + for (const chunk of seeds.data) { + validatePackedChunk(chunk, `${id} seed chunk`); + } +} + +/** Requires one caller-owned packed scalar chunk with its exact output or status length. */ +function validateSingleChunkVector( + vector: GPUVector<'uint32'>, + length: number, + name: string +): void { + validatePackedVector(vector, name); + if (vector.data.length !== 1) { + 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`); + } + validatePackedChunk(vector.data[0], name); + if (vector.data[0].length !== length) { + throw new Error(`${name} chunk must contain exactly ${length} uint32 rows`); + } +} + +/** Requires exact vector-level packed uint32 scalar metadata. */ +function validatePackedVector(vector: GPUVector<'uint32'>, name: string): void { + if ( + 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 packed uint32 rows`); + } +} + +/** Requires exact chunk-level packed uint32 scalar metadata and aligned physical offsets. */ +function validatePackedChunk(chunk: GPUData<'uint32'>, name: string): void { + if ( + chunk.format !== 'uint32' || + !Number.isSafeInteger(chunk.length) || + chunk.length < 0 || + chunk.length > MAXIMUM_UINT32 || + 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 packed, uint32-aligned rows`); + } +} + +/** Keeps every writable result physically disjoint from all graph sources, controls, and peers. */ +function validateDistinctSearchOutputs(search: LuGraphBreadthFirstSearch): void { + const inputVectors = [ + search.topology.graph.sourceVertices, + search.topology.graph.targetVertices, + ...(search.topology.graph.edgeWeights ? [search.topology.graph.edgeWeights] : []), + ...(search.topology.graph.edgeIds ? [search.topology.graph.edgeIds] : []), + ...getAdjacencyVectors(search.topology.forward), + ...(search.topology.reverse ? getAdjacencyVectors(search.topology.reverse) : []), + search.topology.invalidEdgeCount, + search.seeds, + ...(search.seedCount ? [search.seedCount] : []), + ...(search.activeDepth ? [search.activeDepth] : []) + ]; + const physicalAllocations = new Set(); + for (const vector of inputVectors) { + for (const chunk of vector.data) { + physicalAllocations.add(getPhysicalBuffer(chunk)); + } + } + + const outputs = [ + {name: 'distances', vector: search.distances}, + {name: 'predecessors', vector: search.predecessors}, + ...(search.mask ? [{name: 'mask', vector: search.mask}] : []) + ]; + for (const {name, vector} of outputs) { + const physicalBuffer = getPhysicalBuffer(vector.data[0]); + if (physicalAllocations.has(physicalBuffer)) { + throw new Error(`${search.id} ${name} must use a distinct physical buffer allocation`); + } + physicalAllocations.add(physicalBuffer); + } +} + +/** Enumerates existing caller-owned adjacency and status vectors without copying 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 replaceable 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-breadth-first-search.node.spec.ts b/modules/experimental/test/lugraph/lu-graph-breadth-first-search.node.spec.ts new file mode 100644 index 0000000000..5d2141f8b3 --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-breadth-first-search.node.spec.ts @@ -0,0 +1,550 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// 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, + LuGraphBreadthFirstSearch, + LuGraphTopology, + type LuGraphAdjacency, + type LuGraphBreadthFirstSearchDirection, + type LuGraphBreadthFirstSearchProps +} 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 SearchFixture = { + device: NullDevice; + buffers: Buffer[]; + dynamicBuffers: DynamicBuffer[]; + vectors: GPUVector[]; +}; + +type VectorOptions = { + buffer?: Buffer | DynamicBuffer; + byteOffset?: number; + byteStride?: number; + rowByteLength?: number; + stride?: number; +}; + +const searchFixtures: SearchFixture[] = []; + +afterEach(() => { + vi.restoreAllMocks(); + for (const fixture of searchFixtures.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('LuGraphBreadthFirstSearch public API and caller ownership', () => { + test('publishes shortest-path traversal only through the optional luGraph subpath', () => { + expect(typeof LuGraphBreadthFirstSearch).toBe('function'); + expect('LuGraphBreadthFirstSearch' in experimentalModule).toBe(false); + }); + + test('preserves topology, chunked seeds, outputs, and controls without executing GPU work', () => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, { + reverse: true, + weighted: true, + mask: true, + controls: 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 search = new LuGraphBreadthFirstSearch({...props, id: 'borrowed-search'}); + + expect(search.id).toBe('borrowed-search'); + expect(search.topology).toBe(props.topology); + expect(search.seeds).toBe(props.seeds); + expect(search.seedCount).toBe(props.seedCount); + expect(search.distances).toBe(props.distances); + expect(search.predecessors).toBe(props.predecessors); + expect(search.mask).toBe(props.mask); + expect(search.activeDepth).toBe(props.activeDepth); + expect(search.maxDepth).toBe(1); + expect(search.direction).toBe('outgoing'); + expect(search.seeds.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(search, 'destroy')).toBe(false); + + for (const vector of fixture.vectors) vector.destroy(); + expect(fixture.buffers.every(buffer => !buffer.destroyed)).toBe(true); + }); + + test('accepts empty graphs, empty chunked seeds, and optional masks', () => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {vertexCount: 0}); + const seeds = createVector(fixture, 'empty-seeds', 'uint32', []); + const search = new LuGraphBreadthFirstSearch({...props, seeds, maxDepth: 0}); + + expect(search.seeds.data).toEqual([]); + expect(search.distances.length).toBe(0); + expect(search.predecessors.length).toBe(0); + expect(search.mask).toBeUndefined(); + expect(search.maxDepth).toBe(0); + }); +}); + +describe('LuGraphBreadthFirstSearch direction and depth contracts', () => { + test.each([ + 'outgoing', + 'incoming', + 'both' + ] as const)('accepts a directed direction when reverse topology is available: %s', direction => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {reverse: true}); + + expect(new LuGraphBreadthFirstSearch({...props, direction}).direction).toBe(direction); + }); + + test.each([ + 'outgoing', + 'incoming', + 'both' + ] as const)('accepts every undirected direction without redundant reverse topology: %s', direction => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {directed: false}); + + expect(new LuGraphBreadthFirstSearch({...props, direction}).direction).toBe(direction); + }); + + test.each([ + 'incoming', + 'both' + ] as const)('requires reverse adjacency for directed %s traversal', direction => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + + expect(() => new LuGraphBreadthFirstSearch({...props, direction})).toThrow( + /direction|incoming|both|reverse|directed/ + ); + }); + + test.each([ + 'reverse', + 'outbound', + '' + ])('rejects unsupported traversal direction: %s', invalidDirection => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {reverse: true}); + const direction = invalidDirection as LuGraphBreadthFirstSearchDirection; + + expect(() => new LuGraphBreadthFirstSearch({...props, direction})).toThrow(/direction/); + }); + + test.each([0, 1, 1024])('accepts a bounded maximum traversal depth: %i', maxDepth => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + + expect(new LuGraphBreadthFirstSearch({...props, maxDepth}).maxDepth).toBe(maxDepth); + }); + + test.each([ + -1, + 1.5, + Number.NaN, + Number.POSITIVE_INFINITY, + 1025 + ])('rejects an invalid or excessive traversal depth: %s', maxDepth => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + + expect(() => new LuGraphBreadthFirstSearch({...props, maxDepth})).toThrow( + /maxDepth|depth|1024/ + ); + }); +}); + +describe('LuGraphBreadthFirstSearch vector and control validation', () => { + test.each([ + 'distances', + 'predecessors', + 'mask' + ] as const)('requires a packed uint32 vertex output: %s', outputName => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {mask: true}); + const wrongFormat = createVector(fixture, `${outputName}-float`, 'float32', [ + new Float32Array(props.topology.graph.vertexCount) + ]) as unknown as GPUVector<'uint32'>; + + expect(() => new LuGraphBreadthFirstSearch({...props, [outputName]: wrongFormat})).toThrow( + new RegExp(`${outputName}|uint32|packed`) + ); + }); + + test.each([ + 'distances', + 'predecessors', + 'mask' + ] as const)('requires exactly one row per vertex for %s', outputName => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {mask: true}); + const shortOutput = createVector(fixture, `${outputName}-short`, 'uint32', [ + new Uint32Array(props.topology.graph.vertexCount - 1) + ]); + + expect(() => new LuGraphBreadthFirstSearch({...props, [outputName]: shortOutput})).toThrow( + new RegExp(`${outputName}|vertexCount|length`) + ); + }); + + test.each([ + 'distances', + 'predecessors', + 'mask' + ] as const)('requires exactly one physical output chunk for %s', outputName => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {mask: true}); + const partitioned = createVector(fixture, `${outputName}-partitioned`, 'uint32', [ + new Uint32Array(3), + new Uint32Array(3) + ]); + + expect(() => new LuGraphBreadthFirstSearch({...props, [outputName]: partitioned})).toThrow( + new RegExp(`${outputName}|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 predecessor output: %s', (_name, options) => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + const predecessors = createVector( + fixture, + 'unpacked-predecessors', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + options + ); + + expect(() => new LuGraphBreadthFirstSearch({...props, predecessors})).toThrow( + /predecessors|packed|aligned|uint32/ + ); + }); + + test('requires packed uint32 seed chunks and validates intentionally empty chunks', () => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + const wrongFormat = createVector(fixture, 'float-seeds', 'float32', [ + Float32Array.from([0, 1]) + ]) as unknown as GPUVector<'uint32'>; + + expect(() => new LuGraphBreadthFirstSearch({...props, seeds: wrongFormat})).toThrow( + /seeds|uint32|packed/ + ); + Object.defineProperty(props.seeds.data[1], 'format', {value: 'float32'}); + expect(() => new LuGraphBreadthFirstSearch(props)).toThrow(/seeds|uint32|packed|chunk/); + }); + + test.each([ + 'seedCount', + 'activeDepth' + ] as const)('requires one packed uint32 scalar for dynamic %s', controlName => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {controls: true}); + + for (const length of [0, 2]) { + const control = createVector(fixture, `${controlName}-${length}`, 'uint32', [ + new Uint32Array(length) + ]); + expect(() => new LuGraphBreadthFirstSearch({...props, [controlName]: control})).toThrow( + new RegExp(`${controlName}|one|scalar|row`) + ); + } + + const wrongFormat = createVector(fixture, `${controlName}-float`, 'float32', [ + new Float32Array(1) + ]) as unknown as GPUVector<'uint32'>; + expect(() => new LuGraphBreadthFirstSearch({...props, [controlName]: wrongFormat})).toThrow( + new RegExp(`${controlName}|uint32|packed`) + ); + }); + + test('accepts uint32-aligned seed, output, and scalar views at non-256-byte offsets', () => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture); + const seeds = createVector(fixture, 'offset-seeds', 'uint32', [Uint32Array.from([0, 1])], { + byteOffset: 4 + }); + const distances = createVector( + fixture, + 'offset-distances', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + {byteOffset: 4} + ); + const seedCount = createVector(fixture, 'offset-seed-count', 'uint32', [new Uint32Array(1)], { + byteOffset: 4 + }); + + const search = new LuGraphBreadthFirstSearch({...props, seeds, distances, seedCount}); + expect(search.seeds.data[0].byteOffset).toBe(4); + expect(search.distances.data[0].byteOffset).toBe(4); + expect(search.seedCount?.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', + 'seeds', + 'seedCount', + 'activeDepth' + ])('rejects writable distances backed by existing physical allocation: %s', vectorName => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, { + vertexCount: 1, + reverse: true, + weighted: true, + mask: true, + controls: true + }); + const vector = getSearchInputVector(props, vectorName); + const distances = createVector(fixture, 'aliased-distances', 'uint32', [new Uint32Array(1)], { + buffer: vector.data[0].buffer + }); + + expect(() => new LuGraphBreadthFirstSearch({...props, distances})).toThrow( + /distances|distinct|physical|allocation/ + ); + }); + + test.each([ + ['predecessors', 'distances'], + ['mask', 'distances'], + ['mask', 'predecessors'] + ] as const)('rejects aliases between writable %s and %s outputs', (outputName, sourceName) => { + const fixture = createSearchFixture(); + const props = createSearchProps(fixture, {mask: true}); + const sourceVector = props[sourceName]!; + const output = createVector( + fixture, + `aliased-${outputName}`, + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + {buffer: sourceVector.data[0].buffer} + ); + + expect(() => new LuGraphBreadthFirstSearch({...props, [outputName]: output})).toThrow( + /distinct|physical|allocation/ + ); + }); + + test('unwraps borrowed DynamicBuffer objects before checking physical output aliases', () => { + const fixture = createSearchFixture(); + const props = createSearchProps(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 distances = createVector( + fixture, + 'dynamic-aliased-distances', + 'uint32', + [new Uint32Array(props.topology.graph.vertexCount)], + {buffer: dynamicBuffer} + ); + + expect(() => new LuGraphBreadthFirstSearch({...props, distances})).toThrow( + /distinct|physical|allocation/ + ); + expect(concreteBuffer.destroyed).toBe(false); + }); +}); + +function createSearchFixture(): SearchFixture { + const fixture = {device: new NullDevice({}), buffers: [], dynamicBuffers: [], vectors: []}; + searchFixtures.push(fixture); + return fixture; +} + +function createSearchProps( + fixture: SearchFixture, + options: { + vertexCount?: number; + directed?: boolean; + reverse?: boolean; + weighted?: boolean; + mask?: boolean; + controls?: boolean; + } = {} +): LuGraphBreadthFirstSearchProps { + 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 seeds = createVector(fixture, 'seeds', 'uint32', [ + Uint32Array.from([0, 3]), + new Uint32Array(0), + Uint32Array.from([2, 99, 4]) + ]); + const distances = createVector(fixture, 'distances', 'uint32', [new Uint32Array(vertexCount)]); + const predecessors = createVector(fixture, 'predecessors', 'uint32', [ + new Uint32Array(vertexCount) + ]); + const mask = options.mask + ? createVector(fixture, 'mask', 'uint32', [new Uint32Array(vertexCount)]) + : undefined; + const seedCount = options.controls + ? createVector(fixture, 'seedCount', 'uint32', [new Uint32Array(1)]) + : undefined; + const activeDepth = options.controls + ? createVector(fixture, 'activeDepth', 'uint32', [new Uint32Array(1)]) + : undefined; + return {topology, seeds, distances, predecessors, mask, seedCount, activeDepth}; +} + +function createAdjacency( + fixture: SearchFixture, + 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 getSearchInputVector(props: LuGraphBreadthFirstSearchProps, name: string): GPUVector { + if (name === 'seeds') return props.seeds; + if (name === 'seedCount') return props.seedCount!; + if (name === 'activeDepth') return props.activeDepth!; + if (name === 'invalidEdgeCount') return props.topology.invalidEdgeCount; + if (name === 'sourceVertices') return props.topology.graph.sourceVertices; + if (name === 'targetVertices') return props.topology.graph.targetVertices; + if (name === 'edgeWeights') return props.topology.graph.edgeWeights!; + if (name === 'edgeIds') return props.topology.graph.edgeIds!; + + const [direction, vectorName] = name.split('.'); + const adjacency = direction === 'forward' ? props.topology.forward : props.topology.reverse!; + return adjacency[vectorName as keyof LuGraphAdjacency]!; +} + +function createVector( + fixture: SearchFixture, + 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-breadth-first-search.spec.ts b/modules/experimental/test/lugraph/lu-graph-breadth-first-search.spec.ts new file mode 100644 index 0000000000..6b4d647e82 --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-breadth-first-search.spec.ts @@ -0,0 +1,829 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// Copyright (c) vis.gl contributors + +import {Buffer, type Device} from '@luma.gl/core'; +import {GPUCommandGraph} from '@luma.gl/experimental'; +import { + LuGraph, + LuGraphBreadthFirstSearch, + LuGraphTopology, + type LuGraphAdjacency, + type LuGraphBreadthFirstSearchDirection +} 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 { + addLuGraphBreadthFirstSearchToGraphWithDispatchLimit, + getLuGraphBreadthFirstSearchDispatchLayout +} from '../../src/lugraph/lu-graph-breadth-first-search-internals'; + +const UNREACHABLE_VERTEX = 0xffffffff; + +type ScalarFormat = 'uint32' | 'float32'; + +type SearchScenario = { + name: string; + vertexCount: number; + sourceChunks: number[][]; + targetChunks: number[][]; + seedChunks: number[][]; + weightChunks?: number[][]; + directed?: boolean; + reverse?: boolean; + direction?: LuGraphBreadthFirstSearchDirection; + maxDepth?: number; + activeSeedCount?: number; + activeDepth?: number; + mask?: boolean; + capacity?: number; + reverseCapacity?: number; + maximumWorkgroups?: number; + byteOffset?: number; + assertNoScratch?: boolean; +}; + +type ExpectedSearch = { + distances: number[]; + predecessors: number[]; + mask: number[]; + invalidEdgeCount: number; + forwardCount: number; + reverseCount: number; + forwardOverflow: boolean; + reverseOverflow: boolean; + requiredOverflow: boolean; +}; + +type SearchExecutionFixture = { + device: Device; + buffers: Buffer[]; + vectors: GPUVector[]; + graph: LuGraph; + topology: LuGraphTopology; + search: LuGraphBreadthFirstSearch; + commandGraph: GPUCommandGraph; + compiled?: ReturnType; +}; + +const searchScenarios: SearchScenario[] = [ + { + name: 'empty graphs preserve zero-length caller-owned shortest-path outputs', + vertexCount: 0, + sourceChunks: [], + targetChunks: [], + seedChunks: [], + capacity: 0 + }, + { + name: 'isolated vertices retain unreachable sentinels around a valid root', + vertexCount: 6, + sourceChunks: [[], []], + targetChunks: [[], []], + seedChunks: [[2], []], + capacity: 0, + maxDepth: 3 + }, + { + name: 'directed chains publish exact hop distances, predecessors, and source-aligned masks', + vertexCount: 6, + sourceChunks: [[0, 1], [], [2, 3]], + targetChunks: [[1, 2], [], [3, 4]], + seedChunks: [[0]], + maxDepth: 4, + assertNoScratch: true + }, + { + name: 'diamond ties select the lowest predecessor despite shuffled duplicate CSR edges', + vertexCount: 5, + sourceChunks: [[0, 0], [], [2, 1, 1]], + targetChunks: [[2, 1], [], [3, 3, 3]], + seedChunks: [[0]], + maxDepth: 3 + }, + { + name: 'multiple chunked seeds retain roots and deterministic shared-descendant ties', + vertexCount: 6, + sourceChunks: [[4, 1], [], [3, 0]], + targetChunks: [[3, 3], [], [2, 2]], + seedChunks: [[4, 1], [], [99, 4]], + maxDepth: 3 + }, + { + name: 'cycles, duplicate edges, self-loops, and disconnected vertices remain stable', + vertexCount: 7, + sourceChunks: [[0, 1, 1], [], [2, 2, 3]], + targetChunks: [[1, 2, 2], [], [0, 2, 4]], + seedChunks: [[0]], + maxDepth: 6, + mask: false + }, + { + name: 'incoming traversal follows reversed directed CSR and stable predecessor IDs', + vertexCount: 6, + sourceChunks: [[0, 1], [], [1, 4]], + targetChunks: [[2, 2], [], [3, 3]], + seedChunks: [[3]], + reverse: true, + direction: 'incoming', + maxDepth: 3 + }, + { + name: 'bidirectional traversal merges incoming and outgoing deterministic predecessor ties', + vertexCount: 6, + sourceChunks: [[0, 3], [], [2, 4]], + targetChunks: [[2, 0], [], [4, 1]], + seedChunks: [[0]], + reverse: true, + direction: 'both', + maxDepth: 3, + activeDepth: 3 + }, + { + name: 'weighted undirected incoming traversal reuses symmetrized forward adjacency', + vertexCount: 5, + sourceChunks: [[0, 2], [], [2, 3]], + targetChunks: [[1, 1], [], [3, 3]], + weightChunks: [[0.5, 2], [], [4, 8]], + seedChunks: [[3]], + directed: false, + direction: 'incoming', + maxDepth: 4 + }, + { + name: 'undirected bidirectional traversal evaluates forward adjacency only once', + vertexCount: 5, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[2]], + directed: false, + direction: 'both', + maxDepth: 2 + }, + { + name: 'zero maximum depth publishes roots without following graph edges', + vertexCount: 4, + sourceChunks: [[0, 1]], + targetChunks: [[1, 2]], + seedChunks: [[0, 3]], + maxDepth: 0 + }, + { + name: 'bounded traversal leaves vertices beyond the maximum hop unreachable', + vertexCount: 5, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[0]], + maxDepth: 1 + }, + { + name: 'dynamic seed count selects only the leading global rows across seed chunks', + vertexCount: 6, + sourceChunks: [[0, 4, 5]], + targetChunks: [[1, 3, 2]], + seedChunks: [[0, 4], [], [5]], + activeSeedCount: 1, + maxDepth: 2, + activeDepth: 1 + }, + { + name: 'zero dynamic seed count resets every distance, predecessor, and mask', + vertexCount: 4, + sourceChunks: [[0, 1]], + targetChunks: [[1, 2]], + seedChunks: [[0, 2]], + activeSeedCount: 0, + maxDepth: 3, + activeDepth: 2 + }, + { + name: 'dynamic hop counts larger than compiled maximum are safely clamped', + vertexCount: 5, + sourceChunks: [[0, 1, 2, 3]], + targetChunks: [[1, 2, 3, 4]], + seedChunks: [[0]], + maxDepth: 2, + activeDepth: 99 + }, + { + name: 'invalid graph endpoints and invalid seed IDs are ignored without losing valid paths', + vertexCount: 5, + sourceChunks: [[0, 9], [], [2, 3, 4]], + targetChunks: [[1, 2], [], [7, 4, 4]], + seedChunks: [[99, 0], [], [UNREACHABLE_VERTEX, 3]], + maxDepth: 2 + }, + { + name: 'outgoing adjacency overflow fails closed, including otherwise valid roots', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[0]], + capacity: 1, + maxDepth: 3 + }, + { + name: 'incoming adjacency overflow fails closed without publishing partial paths', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[3]], + reverse: true, + reverseCapacity: 0, + direction: 'incoming', + maxDepth: 3 + }, + { + name: 'outgoing traversal ignores overflow from unused reverse adjacency', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[0]], + reverse: true, + reverseCapacity: 0, + direction: 'outgoing', + maxDepth: 3 + }, + { + name: 'incoming traversal ignores overflow from unused forward adjacency', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[3]], + reverse: true, + capacity: 0, + direction: 'incoming', + maxDepth: 3 + }, + { + name: 'bidirectional traversal fails closed when either required adjacency overflows', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[1]], + reverse: true, + reverseCapacity: 2, + direction: 'both', + maxDepth: 2, + activeDepth: 2 + }, + { + name: 'non-256-aligned seed, CSR, output, and dynamic control ranges remain correct', + vertexCount: 5, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + seedChunks: [[0]], + activeSeedCount: 1, + activeDepth: 2, + maxDepth: 3, + byteOffset: 4 + }, + { + name: 'bounded 3D dispatch reaches the final seed and final source vertex of 1025 rows', + vertexCount: 1025, + sourceChunks: [[1024]], + targetChunks: [[512]], + seedChunks: [Array.from({length: 1025}, (_, seedIndex) => (seedIndex === 1024 ? 1024 : 2000))], + maxDepth: 1, + maximumWorkgroups: 2 + } +]; + +test('LuGraphBreadthFirstSearch plans bounded three-dimensional seed and vertex dispatch', tapeTest => { + tapeTest.deepEqual(getLuGraphBreadthFirstSearchDispatchLayout(0, 2), {x: 1, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphBreadthFirstSearchDispatchLayout(512, 2), {x: 2, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphBreadthFirstSearchDispatchLayout(513, 2), {x: 2, y: 2, z: 1}); + tapeTest.deepEqual(getLuGraphBreadthFirstSearchDispatchLayout(1025, 2), {x: 2, y: 2, z: 2}); + tapeTest.throws(() => getLuGraphBreadthFirstSearchDispatchLayout(2049, 2), /3D dispatch limit/); + tapeTest.end(); +}); + +for (const scenario of searchScenarios) { + test(`LuGraphBreadthFirstSearch GPU traversal: ${scenario.name}`, async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const expected = calculateExpectedSearch(scenario); + const fixture = createExecutionFixture(device, scenario, expected); + try { + compileSearch(fixture, scenario.maximumWorkgroups); + executeSearch(fixture); + await assertSearch(tapeTest, fixture, expected); + tapeTest.deepEqual( + fixture.search.seeds.data.map(chunk => chunk.length), + scenario.seedChunks.map(chunk => chunk.length), + 'traversal preserves ordered seed chunks and empty seed batches' + ); + if (scenario.assertNoScratch) { + tapeTest.equal( + fixture.compiled?.stats.logicalTransientBufferCount, + 3, + 'shortest-path traversal adds no frontier or scratch buffers beyond CSR construction' + ); + } + } finally { + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); + }); +} + +test('LuGraphBreadthFirstSearch rereads dynamic seeds, depth, and sources on every encoding', async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const original: SearchScenario = { + name: 'dynamic repeated search', + vertexCount: 6, + sourceChunks: [[0, 1], [], [2, 4]], + targetChunks: [[1, 2], [], [3, 5]], + seedChunks: [[0, 4], [], [5]], + reverse: true, + direction: 'both', + maxDepth: 3, + activeSeedCount: 1, + activeDepth: 1 + }; + const fixture = createExecutionFixture(device, original, calculateExpectedSearch(original)); + const submitSpy = vi.spyOn(device, 'submit'); + const sourceReadbackSpies = [ + ...fixture.graph.sourceVertices.data, + ...fixture.graph.targetVertices.data, + ...fixture.search.seeds.data + ].map(chunk => vi.spyOn(chunk.buffer, 'readAsync')); + + try { + compileSearch(fixture); + tapeTest.equal(submitSpy.mock.calls.length, 0, 'topology and search construction never submit'); + tapeTest.ok( + sourceReadbackSpies.every(spy => spy.mock.calls.length === 0), + 'graph construction and compilation never read source or seed buffers' + ); + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + + executeSearch(fixture); + await assertSearch(tapeTest, fixture, calculateExpectedSearch(original)); + + (fixture.search.seedCount!.data[0].buffer as Buffer).write(Uint32Array.from([2])); + (fixture.search.activeDepth!.data[0].buffer as Buffer).write(Uint32Array.from([3])); + const expanded = {...original, activeSeedCount: 2, activeDepth: 3}; + executeSearch(fixture); + await assertSearch(tapeTest, fixture, calculateExpectedSearch(expanded)); + + const sourceBuffer = fixture.graph.sourceVertices.data[0].buffer as Buffer; + sourceBuffer.write(Uint32Array.from([9, 1])); + const updated = {...expanded, sourceChunks: [[9, 1], [], [2, 4]]}; + executeSearch(fixture); + await assertSearch(tapeTest, fixture, calculateExpectedSearch(updated)); + tapeTest.equal( + fixture.graph.sourceVertices.data[0].buffer, + sourceBuffer, + 'source updates retain the original caller-owned GPUData chunk' + ); + } finally { + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); +}); + +/** Builds complete CPU adjacency before applying deterministic multi-source shortest-hop search. */ +function calculateExpectedSearch(scenario: SearchScenario): ExpectedSearch { + const outgoing = Array.from({length: scenario.vertexCount}, () => [] as number[]); + const incoming = Array.from({length: scenario.vertexCount}, () => [] as number[]); + let invalidEdgeCount = 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; + } + outgoing[source].push(target); + if (scenario.directed === false) { + if (source !== target) outgoing[target].push(source); + } else { + incoming[target].push(source); + } + } + } + + const forwardCount = outgoing.reduce((count, neighbors) => count + neighbors.length, 0); + const reverseCount = incoming.reduce((count, neighbors) => count + neighbors.length, 0); + const forwardOverflow = forwardCount > (scenario.capacity ?? forwardCount); + const reverseOverflow = Boolean( + scenario.reverse && reverseCount > (scenario.reverseCapacity ?? reverseCount) + ); + const direction = scenario.direction ?? 'outgoing'; + const requiredOverflow = + scenario.directed === false || direction === 'outgoing' + ? forwardOverflow + : direction === 'incoming' + ? reverseOverflow + : forwardOverflow || reverseOverflow; + const distances = new Array(scenario.vertexCount).fill(UNREACHABLE_VERTEX); + const predecessors = new Array(scenario.vertexCount).fill(UNREACHABLE_VERTEX); + + if (!requiredOverflow) { + const seeds = scenario.seedChunks.flat(); + const activeSeedCount = Math.min(scenario.activeSeedCount ?? seeds.length, seeds.length); + for (let seedIndex = 0; seedIndex < activeSeedCount; seedIndex++) { + const seed = seeds[seedIndex]; + if (seed < scenario.vertexCount) distances[seed] = 0; + } + + const depth = Math.min(scenario.maxDepth ?? 1, scenario.activeDepth ?? Number.MAX_SAFE_INTEGER); + for (let hop = 0; hop < depth; hop++) { + for (let source = 0; source < scenario.vertexCount; source++) { + if (distances[source] !== hop) continue; + const neighbors = + scenario.directed === false || direction === 'outgoing' + ? outgoing[source] + : direction === 'incoming' + ? incoming[source] + : [...outgoing[source], ...incoming[source]]; + for (const neighbor of neighbors) { + const nextDistance = hop + 1; + if (distances[neighbor] > nextDistance) { + distances[neighbor] = nextDistance; + predecessors[neighbor] = source; + } else if (distances[neighbor] === nextDistance) { + predecessors[neighbor] = Math.min(predecessors[neighbor], source); + } + } + } + } + } + + return { + distances, + predecessors, + mask: distances.map(distance => Number(distance !== UNREACHABLE_VERTEX)), + invalidEdgeCount, + forwardCount, + reverseCount, + forwardOverflow, + reverseOverflow, + requiredOverflow + }; +} + +function createExecutionFixture( + device: Device, + scenario: SearchScenario, + expected: ExpectedSearch +): SearchExecutionFixture { + 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 seeds = createInputVector( + device, + buffers, + vectors, + 'search-seeds', + 'uint32', + scenario.seedChunks, + scenario.byteOffset + ); + const distances = createOutputVector( + device, + buffers, + vectors, + 'search-distances', + 'uint32', + scenario.vertexCount, + scenario.byteOffset + ); + const predecessors = createOutputVector( + device, + buffers, + vectors, + 'search-predecessors', + 'uint32', + scenario.vertexCount, + scenario.byteOffset + ); + const mask = + scenario.mask === false + ? undefined + : createOutputVector( + device, + buffers, + vectors, + 'search-mask', + 'uint32', + scenario.vertexCount, + scenario.byteOffset + ); + const seedCount = + scenario.activeSeedCount === undefined + ? undefined + : createInputVector( + device, + buffers, + vectors, + 'active-seed-count', + 'uint32', + [[scenario.activeSeedCount]], + scenario.byteOffset + ); + const activeDepth = + scenario.activeDepth === undefined + ? undefined + : createInputVector( + device, + buffers, + vectors, + 'active-search-depth', + 'uint32', + [[scenario.activeDepth]], + scenario.byteOffset + ); + const search = new LuGraphBreadthFirstSearch({ + topology, + seeds, + seedCount, + distances, + predecessors, + mask, + maxDepth: scenario.maxDepth, + activeDepth, + direction: scenario.direction + }); + + return { + device, + buffers, + vectors, + graph, + topology, + search, + commandGraph: new GPUCommandGraph(device) + }; +} + +function createInputVector( + device: Device, + buffers: Buffer[], + vectors: GPUVector[], + name: string, + format: Format, + chunks: readonly number[][], + byteOffset = 0 +): 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}`, + byteLength: byteOffset + Math.max(values.length, 1) * Uint32Array.BYTES_PER_ELEMENT, + usage: Buffer.STORAGE | Buffer.COPY_DST + }); + if (values.length > 0) buffer.write(values, byteOffset); + buffers.push(buffer); + return new GPUData({ + buffer, + format, + length: values.length, + byteOffset, + 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 compileSearch(fixture: SearchExecutionFixture, maximumWorkgroups?: number): void { + fixture.topology.addToGraph(fixture.commandGraph); + if (maximumWorkgroups === undefined) { + fixture.search.addToGraph(fixture.commandGraph); + } else { + addLuGraphBreadthFirstSearchToGraphWithDispatchLimit( + fixture.search, + fixture.commandGraph, + maximumWorkgroups + ); + } + fixture.compiled = fixture.commandGraph.compile(); +} + +function executeSearch(fixture: SearchExecutionFixture): void { + const commandEncoder = fixture.device.createCommandEncoder({id: 'lu-graph-breadth-first-test'}); + fixture.compiled!.encode(commandEncoder, {parameters: undefined}); + fixture.device.submit(commandEncoder.finish()); +} + +async function assertSearch( + tapeTest: Test, + fixture: SearchExecutionFixture, + expected: ExpectedSearch +): Promise { + const [distances, predecessors, mask, invalidEdgeCount, forwardOverflow, reverseOverflow] = + await Promise.all([ + readUint32Vector(fixture.search.distances), + readUint32Vector(fixture.search.predecessors), + fixture.search.mask ? readUint32Vector(fixture.search.mask) : Promise.resolve(undefined), + readUint32Vector(fixture.topology.invalidEdgeCount), + readUint32Vector(fixture.topology.forward.overflow), + fixture.topology.reverse + ? readUint32Vector(fixture.topology.reverse.overflow) + : Promise.resolve(undefined) + ]); + + tapeTest.deepEqual(distances, expected.distances, 'shortest-hop distances match the CPU oracle'); + tapeTest.deepEqual( + predecessors, + expected.predecessors, + 'equal-length paths choose the deterministic lowest stable predecessor ID' + ); + if (mask) { + tapeTest.deepEqual(mask, expected.mask, 'optional reachability masks stay source aligned'); + } + tapeTest.equal( + invalidEdgeCount[0], + expected.invalidEdgeCount, + 'invalid graph edges remain excluded' + ); + tapeTest.equal( + forwardOverflow[0], + Number(expected.forwardOverflow), + 'forward overflow remains explicit' + ); + if (reverseOverflow) { + tapeTest.equal( + reverseOverflow[0], + Number(expected.reverseOverflow), + 'reverse overflow remains explicit' + ); + } + if (expected.requiredOverflow) { + tapeTest.ok( + distances.every(distance => distance === UNREACHABLE_VERTEX) && + predecessors.every(predecessor => predecessor === UNREACHABLE_VERTEX), + 'required adjacency overflow fails closed without exposing misleading partial paths' + ); + } +} + +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: SearchExecutionFixture): void { + fixture.compiled?.destroy(); + for (const vector of fixture.vectors) vector.destroy(); + tapeTest.ok( + fixture.buffers.every(buffer => !buffer.destroyed), + 'destroying a compiled traversal and borrowed vectors preserves every caller-owned buffer' + ); + for (const buffer of fixture.buffers) buffer.destroy(); +}