diff --git a/docs/api-reference/experimental/README.md b/docs/api-reference/experimental/README.md index b13c6e3fac..803fefbab0 100644 --- a/docs/api-reference/experimental/README.md +++ b/docs/api-reference/experimental/README.md @@ -114,6 +114,18 @@ schemas, GPU-resident spans, process/thread hierarchy, dependency focus, interac and timeline picking in a dedicated optional submodule. It composes generic command graphs, visibility, flat scenes, and indirect rendering without adding trace concepts to their APIs. +## GPU-resident Graph Analytics + +

+ WebGPU required +

+ +[`@luma.gl/experimental/lugraph`](/docs/api-reference/experimental/lugraph) turns existing GPU +edge columns into reusable compressed adjacency, vertex degrees, bounded shortest-path searches, +weakly connected components, and dangling-aware PageRank scores. Social networks, dependency graphs, +transaction investigations, and infrastructure maps can compose those operations into one WebGPU +command graph without copying source batches or reading complete results back to JavaScript. + ## GPU-resident Linked Crossfiltering

diff --git a/docs/api-reference/experimental/lugraph.md b/docs/api-reference/experimental/lugraph.md new file mode 100644 index 0000000000..4982a6fb17 --- /dev/null +++ b/docs/api-reference/experimental/lugraph.md @@ -0,0 +1,419 @@ +import {ExperimentalDocsTabs} from '@site/src/components/docs/experimental-docs-tabs'; + +# luGraph: GPU-Resident Graph Analytics + + + +## Overview + +A graph answers questions that individual table rows cannot: which accounts share a transaction, +which services depend on a failed service, which people are two introductions apart, and which +pages matter because other important pages link to them. Vertices represent those entities; edges +represent their relationships. + +`@luma.gl/experimental/lugraph` answers these questions directly on a browser WebGPU device. It +describes caller-owned GPU edge columns, builds reusable compressed adjacency, and publishes vertex +degrees, shortest-path neighborhoods, weakly connected groups, and PageRank importance into +caller-owned GPU buffers. Every operation composes with the existing `GPUCommandGraph`. + +This is an experimental, headless graph analytics API, not a graph database, visualization +framework, file importer, or general-purpose dataframe. Applications decide how data reaches the +GPU, which results they render, when commands are submitted, and whether anything is read back. + +## Why keep a graph on the GPU? + +A CPU application can certainly traverse a graph. The problem appears when its relationship data +already lives on the GPU: copying every edge to JavaScript, rebuilding an object graph, running an +analysis, and uploading the answer again interrupts both compute and rendering. + +luGraph keeps the complete intermediate pipeline on one WebGPU device: + +```text +Existing GPU edge columns + -> compressed adjacency + -> degree / shortest paths / weak components / PageRank + -> caller-owned GPU result columns +``` + +The original source and target chunks keep their identities, including empty batches. Adjacency and +analytic outputs remain normal GPU vectors that later compute or application rendering can consume. +Changing a GPU-resident search control and re-encoding an existing compiled graph does not require +materializing a new JavaScript edge list. + +GPU execution is not automatically faster for every graph. A small, CPU-resident, one-off analysis +may be simpler on the CPU because GPU upload, pipeline compilation, command submission, and explicit +readback have real costs. luGraph is most useful when graph data or downstream consumers are already +GPU-resident and several operations reuse the same topology. + +## When should I use luGraph? + +Use luGraph for browser applications that already own typed GPU relationship columns and need to +combine graph analytics with further GPU work: + +- **Social and communication networks:** count contacts, highlight friends within a bounded number + of introductions, group disconnected networks, and rank influential accounts. +- **Software and service dependencies:** follow incoming or outgoing dependency chains, find + isolated dependency islands, and identify packages that many important packages depend on. +- **Transaction and fraud investigations:** follow transfers around a selected account, identify + connected groups of counterparties, and prioritize structurally important entities. +- **Transport and infrastructure maps:** inspect junction degree, unweighted hop reachability, + disconnected subnetworks, and relationship-driven importance across a network. +- **Knowledge and citation graphs:** follow citation links, identify connected collections, and + rank documents by incoming influence rather than raw citation count alone. + +Choose another tool when the application needs weighted shortest paths, a graph query language, +automatic CPU fallback, distributed execution, or compatibility with a CUDA or Python graph API. +luGraph currently operates on one browser WebGPU device and intentionally does not provide those +features. + +## Choose the right graph operation + +| Operation | Question it answers | GPU result | Typical bounded work | +| --- | --- | --- | --- | +| `LuGraph` | Which GPU columns describe the graph? | Borrowed graph metadata and original chunks | Metadata only; no GPU dispatch | +| `LuGraphTopology` | Which vertices are adjacent? | Forward and optional reverse compressed adjacency | `O(V + E)` | +| `LuGraphDegree` | How many relationships touch each vertex in one direction? | One `uint32` degree per vertex | `O(V)` after adjacency exists | +| `LuGraphBreadthFirstSearch` | Which vertices are within a chosen number of unweighted hops? | Distances, deterministic predecessors, and an optional selection mask | At most `O(D × (V + E))` for `D` compiled hops | +| `LuGraphConnectedComponents` | Which vertices belong to the same weakly connected group? | One `uint32` component identifier per vertex | At most `O(K × (V + E))` for `K` bounded iterations | +| `LuGraphPageRank` | Which vertices receive influence from other important vertices? | One normalized `float32` score per vertex | `O(K × (V + E))` for `K` iterations | + +`V` is the graph's explicit vertex count and `E` is its source-edge count. Undirected adjacency +contains both directions for ordinary edges; an undirected self-loop appears once. + +## Describe existing relationships with LuGraph + +**Question: Which existing GPU columns describe the people, accounts, services, or documents in +this network?** + +`LuGraph` is the ownership-preserving entry point. Construct it when an application already has +aligned `GPUVector<'uint32'>` source and target identifiers and knows how many vertices exist, +including isolated vertices that never appear in an edge. + +```ts +import {LuGraph} from '@luma.gl/experimental/lugraph'; + +const graph = new LuGraph({ + vertexCount, + sourceVertices, + targetVertices, + edgeIds, + nodeAttributes, + directed: true +}); +``` + +The graph borrows its vectors; it does not allocate a graph buffer, copy or concatenate chunks, +submit commands, or take ownership of source allocations. Source and target chunks must have the +same ordered lengths. Optional stable edge identifiers and `float32` edge weights follow the same +source partitions, while optional vertex and edge property tables retain their existing metadata. + +Use this lightweight representation when existing GPU tables or render inputs already describe a +network. It is a description, not an upload helper: first create or adapt your GPU vectors through +the application or the appropriate data adapter. + +## Build reusable adjacency with LuGraphTopology + +**Question: Given a particular vertex, which other vertices does it connect to?** + +An edge list answers “what are all relationships?” but repeatedly scanning every edge to discover +one vertex's neighbors is expensive. `LuGraphTopology` builds compressed sparse row (CSR) adjacency +once so later operations can find each vertex's neighbor interval from adjacent offsets. + +For example, a transaction list might contain millions of transfers while an investigator wants +only the accounts directly connected to account 42. Its CSR offset interval identifies that +account's neighbors without asking each later analysis to rescan the entire edge list. + +```ts +import {LuGraphTopology} from '@luma.gl/experimental/lugraph'; + +const topology = new LuGraphTopology({ + graph, + forward: { + offsets: outgoingOffsets, + neighbors: outgoingNeighbors, + edgeIds: outgoingEdgeIds, + count: outgoingCount, + overflow: outgoingOverflow + }, + reverse: { + offsets: incomingOffsets, + neighbors: incomingNeighbors, + edgeIds: incomingEdgeIds, + count: incomingCount, + overflow: incomingOverflow + }, + invalidEdgeCount +}); +``` + +Every shown output is an existing, caller-owned, single-chunk `GPUVector<'uint32'>`. Offsets have +`vertexCount + 1` rows; neighbors and edge identifiers have equal explicit capacities; `count`, +`overflow`, and `invalidEdgeCount` each have one row. When the source graph supplies edge weights, +each configured adjacency also requires a matching `float32` edge-weight output. + +Build reverse adjacency when a directed graph needs incoming-degree queries, incoming or +bidirectional breadth-first search, or PageRank. Directed weak components use forward adjacency +alone. Undirected graphs use one symmetric forward adjacency and must not provide reverse +adjacency. + +Invalid endpoints are excluded and counted. `count` reports the complete number of accepted +adjacency entries even if neighbor capacity is insufficient; `overflow` makes truncation explicit. +Neighbor order within each vertex is intentionally unspecified. + +## Count relationships with LuGraphDegree + +**Question: How many direct relationships does each vertex have?** + +`LuGraphDegree` answers the simplest structural question: how many outgoing or incoming +relationships does each vertex have? Use it to identify network hubs, size junction markers, +detect isolated accounts, or find unusually connected infrastructure and dependency nodes. + +```ts +import {LuGraphDegree} from '@luma.gl/experimental/lugraph'; + +const degree = new LuGraphDegree({ + topology, + output: outgoingDegrees, + direction: 'outgoing' +}); +``` + +Its caller-owned output has one packed `uint32` row per vertex. Outgoing degree is the default; +incoming degree on a directed graph requires reverse adjacency. Duplicate edges count individually, +and an undirected self-loop counts once. + +Degrees come from complete CSR offsets rather than the capacity-bounded neighbor list, so they +remain exact even when the corresponding adjacency reports neighbor overflow. Degree is useful +when raw connectivity is the question; it does not account for whether a vertex's neighbors are +themselves important. + +## Follow unweighted paths with LuGraphBreadthFirstSearch + +**Question: Which entities can I reach within a chosen number of hops, and what shortest path gets +me there?** + +`LuGraphBreadthFirstSearch` expands outward from one or more selected vertices and records the +shortest unweighted hop count to every reachable vertex. Use it to highlight a selected account's +neighborhood, follow a service's dependencies, or explain how two entities connect. + +For example, searching two hops from an account finds both its direct counterparties and the +counterparties of those counterparties. Choose breadth-first search over degree when the question +depends on indirect relationships; choose it over connected components when distance, direction, +or a particular starting vertex matters. + +```ts +import {LuGraphBreadthFirstSearch} from '@luma.gl/experimental/lugraph'; + +const search = new LuGraphBreadthFirstSearch({ + topology, + seeds: selectedVertexIds, + distances: hopDistances, + predecessors: pathParents, + mask: neighborhoodMask, + direction: 'both', + maxDepth: 6, + activeDepth +}); +``` + +`outgoing` follows source-to-target relationships; `incoming` follows their reverse; `both` combines +them. Directed incoming and bidirectional searches require reverse adjacency. `maxDepth` bounds the +number of compiled passes, while an optional one-row GPU `activeDepth` can lower the active search +depth between encodings without rebuilding the command graph. An optional GPU `seedCount` similarly +limits which existing seed rows are active. + +Reached roots have distance zero. Unreachable vertices and root predecessors contain `0xffffffff`; +equal-length parent ties select the numerically lowest stable vertex identifier. Invalid seeds are +ignored, duplicate seeds are harmless, and the optional mask publishes zero or one per vertex. +This is an unweighted shortest-path operation, not a weighted route or travel-time solver. + +## Find disconnected groups with LuGraphConnectedComponents + +**Question: Which vertices belong to the same connected island if edge direction is ignored?** + +`LuGraphConnectedComponents` identifies vertices connected by any path when edge direction is +ignored. Use it to separate disconnected social networks, collect related transaction accounts, +find infrastructure islands, or detect independent dependency groups. + +For example, two transfers `Ana -> Bo` and `Bo -> Cy` put all three accounts in the same group, +even though Cy has no outgoing transfer. An unrelated transfer `Dee -> Eli` forms a different +group. Choose weak components when group membership matters, not the distance from a selected +account or the direction in which influence flows. + +```ts +import {LuGraphConnectedComponents} from '@luma.gl/experimental/lugraph'; + +const components = new LuGraphConnectedComponents({ + topology, + output: componentIds, + iterations: 32, + converged: componentsConverged +}); +``` + +Once propagation converges, every vertex in a weakly connected component receives that group's +lowest stable vertex identifier; an isolated vertex labels itself. Directed edges connect both +endpoints, so reverse adjacency is unnecessary. + +The caller chooses a bounded iteration budget. The optional one-row `uint32` `converged` result is +one only when the final iteration reaches a fixed point; zero means convergence was not established +or the required adjacency overflowed. A connected component answers whether entities connect at +all; it does not claim to discover densely connected communities within one connected network. + +## Rank incoming influence with LuGraphPageRank + +**Question: Which vertices receive influence from other important vertices?** + +`LuGraphPageRank` estimates vertex importance from the importance flowing through incoming +relationships. A citation from an influential paper or a dependency from an important package can +matter more than many links from otherwise disconnected vertices. + +Use PageRank when raw degree is not enough: prioritize influential accounts, rank connected +documents, identify widely depended-on services, or choose salient vertices for application-owned +visualization. The metric is unweighted even when the source topology retains edge-weight columns. + +In a directed graph, `paper A -> paper B` contributes influence from A to B. A paper cited by one +highly influential source can outrank a paper cited by several obscure sources. Degree would count +those citations without asking how influential their sources are; PageRank propagates that +additional context through the surrounding network. + +```ts +import {LuGraphPageRank} from '@luma.gl/experimental/lugraph'; + +const importance = new LuGraphPageRank({ + topology, + output: importanceScores, + damping: 0.85, + iterations: 40, + residual: finalRankChange +}); +``` + +Directed graphs require reverse CSR so each vertex can gather incoming influence; undirected graphs +reuse their symmetric forward adjacency. Every fixed iteration redistributes probability from +dangling vertices with no outgoing edges, applies teleportation, and normalizes the published +`float32` scores so their total is approximately one. + +The default damping is `0.85`: each iteration models an 85% chance of following an outgoing link +and a 15% chance of jumping to a uniformly chosen vertex. This prevents disconnected or cyclic +regions from permanently trapping all influence. A dangling vertex has no outgoing link to follow; +its influence is redistributed uniformly instead of disappearing from the probability vector. +The default bounded iteration count is `40`. + +The optional one-row `float32` `residual` reports the final iteration's L1 score change: the sum of +absolute differences between the last two normalized score vectors. It is an observable error +signal, not an automatic convergence threshold, early-termination mechanism, or promise that a +fixed budget reached the stationary distribution. Reductions use portable WebGPU workgroups and +ordinary `float32` arithmetic, not floating-point atomics or native GPU `float64`. + +## Compose one GPU-resident workflow + +All graph contributors add work to the same caller-owned `GPUCommandGraph`. The following example +assumes that the source columns, packed result vectors, and one-row status vectors already exist +on the same WebGPU device: + +```ts +import {GPUCommandGraph} from '@luma.gl/experimental'; +import { + LuGraph, + LuGraphBreadthFirstSearch, + LuGraphConnectedComponents, + LuGraphDegree, + LuGraphPageRank, + LuGraphTopology +} from '@luma.gl/experimental/lugraph'; + +const graph = new LuGraph({ + vertexCount, + sourceVertices, + targetVertices, + directed: true +}); + +const topology = new LuGraphTopology({ + graph, + forward: { + offsets: outgoingOffsets, + neighbors: outgoingNeighbors, + edgeIds: outgoingEdgeIds, + count: outgoingCount, + overflow: outgoingOverflow + }, + reverse: { + offsets: incomingOffsets, + neighbors: incomingNeighbors, + edgeIds: incomingEdgeIds, + count: incomingCount, + overflow: incomingOverflow + }, + invalidEdgeCount +}); + +const workflow = new GPUCommandGraph(device); + +topology.addToGraph(workflow); +new LuGraphDegree({topology, output: outgoingDegrees}).addToGraph(workflow); +new LuGraphBreadthFirstSearch({ + topology, + seeds: selectedVertexIds, + distances: hopDistances, + predecessors: pathParents, + mask: neighborhoodMask, + direction: 'both', + maxDepth: 6 +}).addToGraph(workflow); +new LuGraphConnectedComponents({ + topology, + output: componentIds, + iterations: 32, + converged: componentsConverged +}).addToGraph(workflow); +new LuGraphPageRank({ + topology, + output: importanceScores, + damping: 0.85, + iterations: 40, + residual: finalRankChange +}).addToGraph(workflow); + +const compiled = workflow.compile(); +const encoder = device.createCommandEncoder({id: 'analyze-network'}); +compiled.encode(encoder, {parameters: undefined}); +device.submit(encoder.finish()); +``` + +Constructors validate existing metadata; they do not upload graph data, submit commands, or read +results. `addToGraph()` declares GPU work, `compile()` resolves the workflow, and the application +explicitly encodes and submits it. Re-encoding rebuilds topology and recomputes the declared +results from the current source and control buffers. + +## Ownership, capacity, and failure boundaries + +- All original source vectors and output vectors are caller-owned. Contributors neither destroy + them nor silently repack their existing chunks. +- Writable outputs require physically distinct GPU buffer allocations, including when a + `DynamicBuffer` wrapper exposes the same underlying allocation through different views. +- Adjacency capacities and overflow statuses are explicit. Breadth-first search fails closed to + unreachable distances, weak components publish `0xffffffff`, and PageRank publishes zero scores + when a required neighbor list overflowed. +- Degree remains exact under neighbor overflow because its input is the complete CSR offset range. +- Fixed component and PageRank iteration budgets do not imply convergence. Their optional status + and final-change outputs remain GPU-resident until an application explicitly requests readback. +- Work uses bounded WebGPU dispatch and portable storage bindings on one device. Original chunk + preservation does not imply distributed or multi-GPU execution. +- The optional graph subpath does not supply automatic Arrow import, rendering, graph persistence, + weighted shortest paths, or a CPU execution fallback. + +See [GPU Primitives and Command Graphs](/docs/api-reference/experimental/gpu-primitives) for the +underlying scheduling, typed GPU vectors, resource ownership, and explicit submission model. + +## Attribution and licensing + +luGraph is inspired by [NVIDIA RAPIDS cuGraph](https://github.com/rapidsai/cugraph) and the NVIDIA +and RAPIDS contributors advancing GPU graph analytics. cuGraph is distributed under the +[Apache License 2.0](https://github.com/rapidsai/cugraph/blob/main/LICENSE). + +This is an independently written, [MIT-licensed](https://github.com/visgl/luma.gl/blob/master/LICENSE) +vis.gl implementation for browser-native WebGPU; it does not copy or translate cuGraph source code. +It does not claim CUDA or cuGraph API compatibility, feature parity, NVIDIA affiliation, or NVIDIA +endorsement. diff --git a/docs/table-of-contents.json b/docs/table-of-contents.json index 37ee37fc20..385c19ecb7 100644 --- a/docs/table-of-contents.json +++ b/docs/table-of-contents.json @@ -209,6 +209,7 @@ "api-reference/experimental/pbr-environment", "api-reference/experimental/geospatial", "api-reference/experimental/luproj", + "api-reference/experimental/lugraph", "api-reference/experimental/luxfilter", "api-reference/experimental/lutrace", "api-reference/experimental/g-buffer", @@ -340,6 +341,7 @@ "api-reference/experimental/pbr-environment", "api-reference/experimental/geospatial", "api-reference/experimental/luproj", + "api-reference/experimental/lugraph", "api-reference/experimental/luxfilter", "api-reference/experimental/lutrace", "api-reference/experimental/g-buffer", diff --git a/modules/experimental/src/lugraph/README.md b/modules/experimental/src/lugraph/README.md index 5b21a20643..e02db6a9d7 100644 --- a/modules/experimental/src/lugraph/README.md +++ b/modules/experimental/src/lugraph/README.md @@ -1,9 +1,16 @@ # @luma.gl/experimental/lugraph -`@luma.gl/experimental/lugraph` provides an optional, headless graph data model over existing, -caller-owned GPU table vectors. Its current foundation preserves source and target vertex columns, -optional edge weights and stable identifiers, property tables, and original chunk boundaries. It -does not upload or copy source data, submit GPU work, render graphs, or provide a graph application. +`@luma.gl/experimental/lugraph` analyzes connected data directly on a browser WebGPU device. Its +optional, headless graph model preserves existing source and target vertex columns, stable edge +identifiers, optional properties, and original GPU vector chunks without uploading or copying them. + +Reusable compressed adjacency supports vertex-degree queries, bounded breadth-first shortest paths, +weakly connected components, and normalized PageRank with dangling-vertex redistribution. Those +operations contribute work to a caller-owned `GPUCommandGraph`; applications retain ownership of +their buffers, rendering, command submission, and any explicitly requested result readback. + +See the [luGraph graph analytics guide](/docs/api-reference/experimental/lugraph) for when to use +each operation, complete GPU-resident composition examples, and ownership and capacity contracts. ## Attribution and licensing diff --git a/modules/experimental/src/lugraph/index.ts b/modules/experimental/src/lugraph/index.ts index 2d1c2a71c6..de09f77600 100644 --- a/modules/experimental/src/lugraph/index.ts +++ b/modules/experimental/src/lugraph/index.ts @@ -16,3 +16,5 @@ export type { } from './lu-graph-breadth-first-search'; export {LuGraphConnectedComponents} from './lu-graph-connected-components'; export type {LuGraphConnectedComponentsProps} from './lu-graph-connected-components'; +export {LuGraphPageRank} from './lu-graph-page-rank'; +export type {LuGraphPageRankProps} from './lu-graph-page-rank'; diff --git a/modules/experimental/src/lugraph/lu-graph-page-rank-internals.ts b/modules/experimental/src/lugraph/lu-graph-page-rank-internals.ts new file mode 100644 index 0000000000..e230daf6d2 --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-page-rank-internals.ts @@ -0,0 +1,603 @@ +// 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 { + createTransientView, + getViewBinding, + getViewElementOffset +} from '../gpu-primitives/graph-data-view-utils'; +import type {LuGraphPageRank} from './lu-graph-page-rank'; + +const PAGE_RANK_WORKGROUP_SIZE = 256; + +type PageRankDataView = GraphDataView<'uint32'> | GraphDataView<'float32'>; + +type ImportedPageRank = { + id: string; + vertexCount: number; + damping: number; + forwardOffsets: GraphDataView<'uint32'>; + incomingOffsets: GraphDataView<'uint32'>; + incomingNeighbors: GraphDataView<'uint32'>; + overflow: GraphDataView<'uint32'>; + reverseOverflow?: GraphDataView<'uint32'>; + output: GraphDataView<'float32'>; + residual?: GraphDataView<'float32'>; + maxComputeWorkgroupsPerDimension: number; +}; + +type PageRankBinding = { + view: PageRankDataView; + usage: GraphBufferUse['usage']; +}; + +type PageRankPassProps = { + id: string; + source: string; + bindings: Record; + dispatchLayout: GPUBoundedDispatchLayout; +}; + +/** Adds dangling-safe GPU PageRank using an explicit bounded dispatch limit. @internal */ +export function addLuGraphPageRankToGraphWithDispatchLimit( + pageRank: LuGraphPageRank, + commandGraph: GPUCommandGraph, + maxComputeWorkgroupsPerDimension: number +): void { + if (pageRank.topology.graph.vertexCount === 0 && !pageRank.residual) { + return; + } + + const directed = pageRank.topology.graph.directed; + const forwardOffsets = commandGraph.importGPUVector( + `${pageRank.id}-forward-offsets`, + pageRank.topology.forward.offsets + ).data[0]; + const incoming = directed ? pageRank.topology.reverse! : pageRank.topology.forward; + const state: ImportedPageRank = { + id: pageRank.id, + vertexCount: pageRank.topology.graph.vertexCount, + damping: pageRank.damping, + forwardOffsets, + incomingOffsets: directed + ? commandGraph.importGPUVector(`${pageRank.id}-incoming-offsets`, incoming.offsets).data[0] + : forwardOffsets, + incomingNeighbors: commandGraph.importGPUVector( + `${pageRank.id}-incoming-neighbors`, + incoming.neighbors + ).data[0], + overflow: commandGraph.importGPUVector( + `${pageRank.id}-forward-overflow`, + pageRank.topology.forward.overflow + ).data[0], + ...(directed + ? { + reverseOverflow: commandGraph.importGPUVector( + `${pageRank.id}-incoming-overflow`, + incoming.overflow + ).data[0] + } + : {}), + output: commandGraph.importGPUVector(`${pageRank.id}-output`, pageRank.output).data[0], + ...(pageRank.residual + ? { + residual: commandGraph.importGPUVector(`${pageRank.id}-residual`, pageRank.residual) + .data[0] + } + : {}), + maxComputeWorkgroupsPerDimension + }; + + addInitializationPass(commandGraph, state); + if (state.vertexCount === 0) { + return; + } + + const workspace = createTransientView( + commandGraph, + `${state.id}-workspace`, + 'float32', + state.vertexCount + ); + const reductionLevels = createReductionLevels(commandGraph, state); + + for (let iteration = 0; iteration < pageRank.iterations; iteration++) { + addDanglingGatherPass(commandGraph, {state, workspace, iteration}); + const danglingMass = addReduction(commandGraph, { + id: `${state.id}-iteration-${iteration}-dangling`, + state, + input: workspace, + levels: reductionLevels + }); + addPullPass(commandGraph, {state, workspace, danglingMass, iteration}); + const rankSum = addReduction(commandGraph, { + id: `${state.id}-iteration-${iteration}-sum`, + state, + input: workspace, + levels: reductionLevels + }); + const collectResidual = Boolean(state.residual && iteration === pageRank.iterations - 1); + addNormalizationPass(commandGraph, {state, workspace, rankSum, iteration, collectResidual}); + if (collectResidual) { + addReduction(commandGraph, { + id: `${state.id}-residual`, + state, + input: workspace, + levels: reductionLevels, + output: state.residual! + }); + } + } +} + +/** Initializes uniform scores and the optional residual while failing closed on overflow. */ +function addInitializationPass( + commandGraph: GPUCommandGraph, + state: ImportedPageRank +): void { + const bindings: Record = { + output: {view: state.output, usage: 'storage-write'}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.reverseOverflow + ? {reverseOverflow: {view: state.reverseOverflow, usage: 'storage-read'}} + : {}), + ...(state.residual ? {residual: {view: state.residual, usage: 'storage-write'}} : {}) + }; + const reverseOffset = state.reverseOverflow + ? `const REVERSE_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.reverseOverflow)}u;` + : ''; + const residualOffset = state.residual + ? `const RESIDUAL_OFFSET: u32 = ${getViewElementOffset(state.residual)}u;` + : ''; + const reverseOverflow = state.reverseOverflow + ? ' || reverseOverflow[REVERSE_OVERFLOW_OFFSET] != 0u' + : ''; + const clearResidual = state.residual + ? 'if (index == 0u) { residual[RESIDUAL_OFFSET] = 0.0; }' + : ''; + const dispatchLayout = getLuGraphPageRankDispatchLayout( + 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; +${reverseOffset} +${residualOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${PAGE_RANK_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, PAGE_RANK_WORKGROUP_SIZE)} + if (index < VERTEX_COUNT) { + let hasOverflow = overflow[OVERFLOW_OFFSET] != 0u${reverseOverflow}; + let uniformScore = 1.0 / f32(max(VERTEX_COUNT, 1u)); + output[OUTPUT_OFFSET + index] = select(uniformScore, 0.0, hasOverflow); + } + ${clearResidual} +}`; + + addPageRankPass(commandGraph, { + id: `${state.id}-initialize`, + source, + bindings, + dispatchLayout + }); +} + +/** Allocates one reusable 256-way reduction hierarchy shared by all ranking iterations. */ +function createReductionLevels( + commandGraph: GPUCommandGraph, + state: ImportedPageRank +): GraphDataView<'float32'>[] { + const levels: GraphDataView<'float32'>[] = []; + let length = state.vertexCount; + do { + length = Math.ceil(length / PAGE_RANK_WORKGROUP_SIZE); + levels.push( + createTransientView( + commandGraph, + `${state.id}-reduction-level-${levels.length}`, + 'float32', + length + ) + ); + } while (length > 1); + return levels; +} + +/** Extracts dangling-node probability mass without unsupported floating-point atomics. */ +function addDanglingGatherPass( + commandGraph: GPUCommandGraph, + props: { + state: ImportedPageRank; + workspace: GraphDataView<'float32'>; + iteration: number; + } +): void { + const {state, workspace} = props; + const bindings: Record = { + output: {view: state.output, usage: 'storage-read'}, + forwardOffsets: {view: state.forwardOffsets, usage: 'storage-read'}, + workspace: {view: workspace, usage: 'storage-write'}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.reverseOverflow + ? {reverseOverflow: {view: state.reverseOverflow, usage: 'storage-read'}} + : {}) + }; + const reverseOffset = state.reverseOverflow + ? `const REVERSE_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.reverseOverflow)}u;` + : ''; + const reverseOverflow = state.reverseOverflow + ? ' || reverseOverflow[REVERSE_OVERFLOW_OFFSET] != 0u' + : ''; + const dispatchLayout = getLuGraphPageRankDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const FORWARD_OFFSETS_OFFSET: u32 = ${getViewElementOffset(state.forwardOffsets)}u; +const WORKSPACE_OFFSET: u32 = ${getViewElementOffset(workspace)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${reverseOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${PAGE_RANK_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, PAGE_RANK_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT) { return; } + let hasOverflow = overflow[OVERFLOW_OFFSET] != 0u${reverseOverflow}; + var contribution = 0.0; + if (!hasOverflow) { + let degree = + forwardOffsets[FORWARD_OFFSETS_OFFSET + index + 1u] - + forwardOffsets[FORWARD_OFFSETS_OFFSET + index]; + if (degree == 0u) { contribution = output[OUTPUT_OFFSET + index]; } + } + workspace[WORKSPACE_OFFSET + index] = contribution; +}`; + + addPageRankPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-gather-dangling`, + source, + bindings, + dispatchLayout + }); +} + +/** Applies the reverse-CSR pull recurrence using no more than eight storage bindings. */ +function addPullPass( + commandGraph: GPUCommandGraph, + props: { + state: ImportedPageRank; + workspace: GraphDataView<'float32'>; + danglingMass: GraphDataView<'float32'>; + iteration: number; + } +): void { + const {state, workspace, danglingMass} = props; + const bindings: Record = { + output: {view: state.output, usage: 'storage-read'}, + workspace: {view: workspace, usage: 'storage-write'}, + forwardOffsets: {view: state.forwardOffsets, usage: 'storage-read'}, + ...(state.reverseOverflow + ? {incomingOffsets: {view: state.incomingOffsets, usage: 'storage-read'}} + : {}), + incomingNeighbors: {view: state.incomingNeighbors, usage: 'storage-read'}, + danglingMass: {view: danglingMass, usage: 'storage-read'}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.reverseOverflow + ? {reverseOverflow: {view: state.reverseOverflow, usage: 'storage-read'}} + : {}) + }; + const incomingOffset = state.reverseOverflow + ? `const INCOMING_OFFSETS_OFFSET: u32 = ${getViewElementOffset(state.incomingOffsets)}u;` + : ''; + const reverseOffset = state.reverseOverflow + ? `const REVERSE_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.reverseOverflow)}u;` + : ''; + const reverseOverflow = state.reverseOverflow + ? ' || reverseOverflow[REVERSE_OVERFLOW_OFFSET] != 0u' + : ''; + const incomingOffsets = state.reverseOverflow ? 'incomingOffsets' : 'forwardOffsets'; + const incomingOffsetsOffset = state.reverseOverflow + ? 'INCOMING_OFFSETS_OFFSET' + : 'FORWARD_OFFSETS_OFFSET'; + const dispatchLayout = getLuGraphPageRankDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const CAPACITY: u32 = ${state.incomingNeighbors.length}u; +const DAMPING: f32 = ${state.damping}; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const WORKSPACE_OFFSET: u32 = ${getViewElementOffset(workspace)}u; +const FORWARD_OFFSETS_OFFSET: u32 = ${getViewElementOffset(state.forwardOffsets)}u; +const INCOMING_NEIGHBORS_OFFSET: u32 = ${getViewElementOffset(state.incomingNeighbors)}u; +const DANGLING_MASS_OFFSET: u32 = ${getViewElementOffset(danglingMass)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${incomingOffset} +${reverseOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${PAGE_RANK_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, PAGE_RANK_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT) { return; } + let hasOverflow = overflow[OVERFLOW_OFFSET] != 0u${reverseOverflow}; + if (hasOverflow) { + workspace[WORKSPACE_OFFSET + index] = 0.0; + return; + } + let first = min(${incomingOffsets}[${incomingOffsetsOffset} + index], CAPACITY); + let last = min(${incomingOffsets}[${incomingOffsetsOffset} + index + 1u], CAPACITY); + var incomingMass = 0.0; + for (var slot = first; slot < last; slot++) { + let neighbor = incomingNeighbors[INCOMING_NEIGHBORS_OFFSET + slot]; + if (neighbor >= VERTEX_COUNT) { continue; } + let degree = + forwardOffsets[FORWARD_OFFSETS_OFFSET + neighbor + 1u] - + forwardOffsets[FORWARD_OFFSETS_OFFSET + neighbor]; + if (degree > 0u) { + incomingMass += output[OUTPUT_OFFSET + neighbor] / f32(degree); + } + } + let vertexCount = f32(VERTEX_COUNT); + let redistributedDangling = danglingMass[DANGLING_MASS_OFFSET] / vertexCount; + let teleportation = (1.0 - DAMPING) / vertexCount; + workspace[WORKSPACE_OFFSET + index] = + teleportation + DAMPING * (incomingMass + redistributedDangling); +}`; + + addPageRankPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-pull`, + source, + bindings, + dispatchLayout + }); +} + +/** Normalizes every iteration and optionally writes final absolute residual contributions. */ +function addNormalizationPass( + commandGraph: GPUCommandGraph, + props: { + state: ImportedPageRank; + workspace: GraphDataView<'float32'>; + rankSum: GraphDataView<'float32'>; + iteration: number; + collectResidual: boolean; + } +): void { + const {state, workspace, rankSum} = props; + const bindings: Record = { + output: {view: state.output, usage: 'storage-read-write'}, + workspace: { + view: workspace, + usage: props.collectResidual ? 'storage-read-write' : 'storage-read' + }, + rankSum: {view: rankSum, usage: 'storage-read'}, + overflow: {view: state.overflow, usage: 'storage-read'}, + ...(state.reverseOverflow + ? {reverseOverflow: {view: state.reverseOverflow, usage: 'storage-read'}} + : {}) + }; + const reverseOffset = state.reverseOverflow + ? `const REVERSE_OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.reverseOverflow)}u;` + : ''; + const reverseOverflow = state.reverseOverflow + ? ' || reverseOverflow[REVERSE_OVERFLOW_OFFSET] != 0u' + : ''; + const collectResidual = props.collectResidual + ? 'workspace[WORKSPACE_OFFSET + index] = difference;' + : ''; + const dispatchLayout = getLuGraphPageRankDispatchLayout( + state.vertexCount, + state.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const VERTEX_COUNT: u32 = ${state.vertexCount}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(state.output)}u; +const WORKSPACE_OFFSET: u32 = ${getViewElementOffset(workspace)}u; +const RANK_SUM_OFFSET: u32 = ${getViewElementOffset(rankSum)}u; +const OVERFLOW_OFFSET: u32 = ${getViewElementOffset(state.overflow)}u; +${reverseOffset} +${getBindingDeclarations(bindings)} + +@compute @workgroup_size(${PAGE_RANK_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, PAGE_RANK_WORKGROUP_SIZE)} + if (index >= VERTEX_COUNT) { return; } + let hasOverflow = overflow[OVERFLOW_OFFSET] != 0u${reverseOverflow}; + let total = rankSum[RANK_SUM_OFFSET]; + let validTotal = total > 0.0 && total == total && abs(total) <= 3.402823466e+38; + var next = 0.0; + var difference = 0.0; + if (!hasOverflow && validTotal) { + next = workspace[WORKSPACE_OFFSET + index] / total; + difference = abs(next - output[OUTPUT_OFFSET + index]); + } + output[OUTPUT_OFFSET + index] = next; + ${collectResidual} +}`; + + addPageRankPass(commandGraph, { + id: `${state.id}-iteration-${props.iteration}-normalize`, + source, + bindings, + dispatchLayout + }); +} + +/** Reuses one bounded workgroup hierarchy for dangling mass, rank sums, and final residual. */ +function addReduction( + commandGraph: GPUCommandGraph, + props: { + id: string; + state: ImportedPageRank; + input: GraphDataView<'float32'>; + levels: GraphDataView<'float32'>[]; + output?: GraphDataView<'float32'>; + } +): GraphDataView<'float32'> { + let input = props.input; + for (const [levelIndex, level] of props.levels.entries()) { + const last = levelIndex === props.levels.length - 1; + const output = last && props.output ? props.output : level; + addReductionPass(commandGraph, { + id: `${props.id}-level-${levelIndex}`, + input, + output, + maxComputeWorkgroupsPerDimension: props.state.maxComputeWorkgroupsPerDimension + }); + input = output; + } + return input; +} + +/** Sums 256 float32 lanes with only workgroup-uniform exits before synchronization barriers. */ +function addReductionPass( + commandGraph: GPUCommandGraph, + props: { + id: string; + input: GraphDataView<'float32'>; + output: GraphDataView<'float32'>; + maxComputeWorkgroupsPerDimension: number; + } +): void { + const bindings: Record = { + inputValues: {view: props.input, usage: 'storage-read'}, + outputValues: {view: props.output, usage: 'storage-write'} + }; + const dispatchLayout = getLuGraphPageRankDispatchLayout( + props.input.length, + props.maxComputeWorkgroupsPerDimension + ); + const source = /* wgsl */ ` +const INPUT_COUNT: u32 = ${props.input.length}u; +const OUTPUT_COUNT: u32 = ${props.output.length}u; +const INPUT_OFFSET: u32 = ${getViewElementOffset(props.input)}u; +const OUTPUT_OFFSET: u32 = ${getViewElementOffset(props.output)}u; +${getBindingDeclarations(bindings)} +var reductionValues: array; + +@compute @workgroup_size(${PAGE_RANK_WORKGROUP_SIZE}) +fn main( + @builtin(workgroup_id) workgroupId: vec3, + @builtin(local_invocation_index) localInvocationIndex: u32 +) { + ${getBoundedInvocationIndexSource(dispatchLayout, PAGE_RANK_WORKGROUP_SIZE)} + if (workgroupIndex >= OUTPUT_COUNT) { return; } + var value = 0.0; + if (index < INPUT_COUNT) { value = inputValues[INPUT_OFFSET + index]; } + reductionValues[localInvocationIndex] = value; + workgroupBarrier(); + + for (var stride = ${PAGE_RANK_WORKGROUP_SIZE / 2}u; stride > 0u; stride /= 2u) { + if (localInvocationIndex < stride) { + reductionValues[localInvocationIndex] += reductionValues[localInvocationIndex + stride]; + } + workgroupBarrier(); + } + if (localInvocationIndex == 0u) { + outputValues[OUTPUT_OFFSET + workgroupIndex] = reductionValues[0]; + } +}`; + + addPageRankPass(commandGraph, {id: props.id, source, bindings, dispatchLayout}); +} + +/** Declares packed uint32 and float32 storage views in generated binding-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.view.format === 'float32' ? 'f32' : 'u32'; + return `@group(0) @binding(${location}) var ${name}: array<${element}>;`; + }) + .join('\n'); +} + +/** Compiles one bounded GPU pass without hidden submission, synchronization, or readback. */ +function addPageRankPass( + commandGraph: GPUCommandGraph, + props: PageRankPassProps +): 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 bounded three-dimensional PageRank vertex and hierarchical-reduction dispatch. @internal */ +export function getLuGraphPageRankDispatchLayout( + elementCount: number, + maxComputeWorkgroupsPerDimension: number +): GPUBoundedDispatchLayout { + return getBoundedDispatchLayout( + 'LuGraphPageRank', + elementCount, + PAGE_RANK_WORKGROUP_SIZE, + maxComputeWorkgroupsPerDimension + ); +} diff --git a/modules/experimental/src/lugraph/lu-graph-page-rank.ts b/modules/experimental/src/lugraph/lu-graph-page-rank.ts new file mode 100644 index 0000000000..61c8662484 --- /dev/null +++ b/modules/experimental/src/lugraph/lu-graph-page-rank.ts @@ -0,0 +1,178 @@ +// 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 {addLuGraphPageRankToGraphWithDispatchLimit} from './lu-graph-page-rank-internals'; +import type {LuGraphAdjacency, LuGraphTopology} from './lu-graph-topology'; + +const DEFAULT_PAGE_RANK_DAMPING = 0.85; +const DEFAULT_PAGE_RANK_ITERATIONS = 40; +const MAXIMUM_PAGE_RANK_ITERATIONS = 1024; +const SCALAR_BYTE_LENGTH = 4; + +/** Existing graph topology, caller-owned PageRank scores, and optional residual. */ +export type LuGraphPageRankProps = { + /** Prefix for generated command-graph nodes and imported resources. */ + id?: string; + /** Existing GPU-resident graph topology; directed graphs require reverse adjacency. */ + topology: LuGraphTopology; + /** One caller-owned, packed floating-point PageRank score for each graph vertex. */ + output: GPUVector<'float32'>; + /** Probability of following an outgoing edge rather than teleporting. Defaults to 0.85. */ + damping?: number; + /** Bounded number of compiled, normalized PageRank iterations. Defaults to 40. */ + iterations?: number; + /** Optional caller-owned scalar receiving the final iteration's absolute rank change. */ + residual?: GPUVector<'float32'>; +}; + +/** + * Publishes normalized, unweighted PageRank scores entirely from existing GPU graph topology. + * + * Directed graphs require reverse adjacency for incoming-edge gathers; undirected graphs reuse + * their symmetric forward adjacency. Each iteration redistributes dangling-vertex mass before + * normalizing the published scores. Existing edge weights do not affect this unweighted metric. + * Overflow in either required adjacency instead publishes zero scores and a zero residual. + */ +export class LuGraphPageRank { + /** Prefix for generated command-graph nodes and imported resources. */ + readonly id: string; + /** Existing caller-owned GPU graph topology. */ + readonly topology: LuGraphTopology; + /** Caller-owned, vertex-aligned floating-point PageRank scores. */ + readonly output: GPUVector<'float32'>; + /** Probability of following an outgoing edge rather than teleporting. */ + readonly damping: number; + /** Number of compiled, synchronized PageRank iterations. */ + readonly iterations: number; + /** Optional caller-owned GPU-resident final absolute rank-change scalar. */ + readonly residual?: GPUVector<'float32'>; + + /** Validates existing caller-owned metadata without allocating, submitting, or reading work. */ + constructor(props: LuGraphPageRankProps) { + this.id = props.id ?? 'lu-graph-page-rank'; + this.topology = props.topology; + this.output = props.output; + this.damping = props.damping ?? DEFAULT_PAGE_RANK_DAMPING; + this.iterations = props.iterations ?? DEFAULT_PAGE_RANK_ITERATIONS; + this.residual = props.residual; + + if (this.topology.graph.directed && !this.topology.reverse) { + throw new Error(`${this.id} directed PageRank requires reverse adjacency`); + } + if (!Number.isFinite(this.damping) || this.damping < 0 || this.damping > 1) { + throw new Error(`${this.id} damping must be a finite number between zero and one`); + } + if ( + !Number.isSafeInteger(this.iterations) || + this.iterations < 1 || + this.iterations > MAXIMUM_PAGE_RANK_ITERATIONS + ) { + throw new Error(`${this.id} iterations must be a safe integer between one and 1024`); + } + + validatePageRankVector(this.output, this.topology.graph.vertexCount, `${this.id} output`); + if (this.residual) { + validatePageRankVector(this.residual, 1, `${this.id} residual`); + } + validateDistinctPageRankOutputs(this); + } + + /** Declares bounded graph ranking work without submitting commands or reading results. */ + addToGraph(commandGraph: GPUCommandGraph): void { + addLuGraphPageRankToGraphWithDispatchLimit( + this, + commandGraph, + commandGraph.device.limits.maxComputeWorkgroupsPerDimension + ); + } +} + +/** Requires one packed, aligned floating-point output chunk with its exact logical row count. */ +function validatePageRankVector(vector: GPUVector<'float32'>, length: number, name: string): void { + if ( + vector.data.length !== 1 || + vector.format !== 'float32' || + 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 float32 chunk`); + } + if (vector.length !== length) { + throw new Error(`${name} must contain exactly ${length} float32 rows`); + } + + const chunk = vector.data[0]; + if ( + chunk.format !== 'float32' || + 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, float32-aligned chunk`); + } +} + +/** Keeps ranking scores and optional residual disjoint from every existing graph allocation. */ +function validateDistinctPageRankOutputs(pageRank: LuGraphPageRank): void { + const topology = pageRank.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: pageRank.output}, + ...(pageRank.residual ? [{name: 'residual', vector: pageRank.residual}] : []) + ]; + for (const {name, vector} of outputs) { + const buffer = getPhysicalBuffer(vector.data[0]); + if (allocations.has(buffer)) { + throw new Error(`${pageRank.id} ${name} must use a distinct physical buffer allocation`); + } + allocations.add(buffer); + } +} + +/** Enumerates existing adjacency and status columns without changing any chunk identities. */ +function getAdjacencyVectors( + adjacency: LuGraphAdjacency +): (GPUVector<'uint32'> | GPUVector<'float32'>)[] { + return [ + adjacency.offsets, + adjacency.neighbors, + adjacency.edgeIds, + ...(adjacency.edgeWeights ? [adjacency.edgeWeights] : []), + adjacency.count, + adjacency.overflow + ]; +} + +/** Resolves an engine wrapper to its current underlying 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-page-rank.node.spec.ts b/modules/experimental/test/lugraph/lu-graph-page-rank.node.spec.ts new file mode 100644 index 0000000000..fe5202b873 --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-page-rank.node.spec.ts @@ -0,0 +1,467 @@ +// 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, + LuGraphPageRank, + LuGraphTopology, + type LuGraphAdjacency, + type LuGraphPageRankProps +} 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 PageRankFixture = { + device: NullDevice; + buffers: Buffer[]; + dynamicBuffers: DynamicBuffer[]; + vectors: GPUVector[]; +}; + +type VectorOptions = { + buffer?: Buffer | DynamicBuffer; + byteOffset?: number; + byteStride?: number; + rowByteLength?: number; + stride?: number; +}; + +const pageRankFixtures: PageRankFixture[] = []; + +afterEach(() => { + vi.restoreAllMocks(); + for (const fixture of pageRankFixtures.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('LuGraphPageRank optional API and caller-owned resources', () => { + test('exposes PageRank only through the experimental luGraph package subpath', () => { + expect(typeof LuGraphPageRank).toBe('function'); + expect('LuGraphPageRank' in experimentalModule).toBe(false); + }); + + test('retains caller-owned topology and float outputs without allocating or executing GPU work', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {weighted: true, residual: 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 pageRank = new LuGraphPageRank({...props, id: 'borrowed-page-rank'}); + + expect(pageRank.id).toBe('borrowed-page-rank'); + expect(pageRank.topology).toBe(props.topology); + expect(pageRank.output).toBe(props.output); + expect(pageRank.residual).toBe(props.residual); + expect(pageRank.damping).toBe(0.85); + expect(pageRank.iterations).toBe(40); + expect(pageRank.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(pageRank, 'destroy')).toBe(false); + + for (const vector of fixture.vectors) vector.destroy(); + expect(fixture.buffers.every(buffer => !buffer.destroyed)).toBe(true); + }); + + test('requires reverse CSR for directed graphs while allowing symmetric undirected topology', () => { + const fixture = createPageRankFixture(); + const directed = createPageRankProps(fixture, {reverse: false}); + const undirected = createPageRankProps(fixture, {directed: false, reverse: false}); + + expect(() => new LuGraphPageRank(directed)).toThrow(/reverse|incoming|directed/); + expect(new LuGraphPageRank(undirected).topology.reverse).toBeUndefined(); + }); + + test('accepts empty rank output with an optional float32 residual scalar', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {vertexCount: 0, residual: true}); + const pageRank = new LuGraphPageRank(props); + + expect(pageRank.output.length).toBe(0); + expect(pageRank.output.data).toHaveLength(1); + expect(pageRank.output.data[0].buffer.byteLength).toBeGreaterThanOrEqual(4); + expect(pageRank.residual?.length).toBe(1); + }); +}); + +describe('LuGraphPageRank bounded parameters and float32 vector validation', () => { + test.each([ + 0, 0.5, 0.85, 1 + ])('accepts a finite damping factor in the closed unit interval: %s', damping => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + + expect(new LuGraphPageRank({...props, damping}).damping).toBe(damping); + }); + + test.each([ + -0.001, + 1.001, + Number.NaN, + Number.POSITIVE_INFINITY, + Number.NEGATIVE_INFINITY + ])('rejects an invalid damping factor: %s', damping => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + + expect(() => new LuGraphPageRank({...props, damping})).toThrow(/damping|finite|between/); + }); + + test.each([1, 40, 1024])('accepts a positive bounded iteration count: %i', iterations => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + + expect(new LuGraphPageRank({...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 = createPageRankFixture(); + const props = createPageRankProps(fixture); + + expect(() => new LuGraphPageRank({...props, iterations})).toThrow(/iterations|positive|1024/); + }); + + test.each([5, 7])('requires exactly one float score per graph vertex: %i', length => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + const output = createVector(fixture, `score-length-${length}`, 'float32', [ + new Float32Array(length) + ]); + + expect(() => new LuGraphPageRank({...props, output})).toThrow(/output|vertexCount|length/); + }); + + test('requires float32 PageRank scores instead of uint32', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + const output = createVector(fixture, 'uint-page-rank', 'uint32', [ + new Uint32Array(props.topology.graph.vertexCount) + ]) as unknown as GPUVector<'float32'>; + + expect(() => new LuGraphPageRank({...props, output})).toThrow(/output|float32|packed/); + }); + + test.each([0, 2])('requires exactly one physical score chunk: %i', chunkCount => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + const chunks = chunkCount === 0 ? [] : [new Float32Array(3), new Float32Array(3)]; + const output = createVector(fixture, 'partitioned-scores', 'float32', chunks); + + expect(() => new LuGraphPageRank({...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 score output: %s', (_name, options) => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + const output = createVector( + fixture, + 'unpacked-page-rank', + 'float32', + [new Float32Array(props.topology.graph.vertexCount)], + options + ); + + expect(() => new LuGraphPageRank({...props, output})).toThrow(/output|packed|aligned|float32/); + }); + + test.each([0, 2])('requires exactly one final float32 residual row: %i', length => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {residual: true}); + const residual = createVector(fixture, `residual-length-${length}`, 'float32', [ + new Float32Array(length) + ]); + + expect(() => new LuGraphPageRank({...props, residual})).toThrow(/residual|one|row|scalar/); + }); + + test('requires packed float32 residual data with exactly one chunk', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {residual: true}); + const wrongFormat = createVector(fixture, 'uint-residual', 'uint32', [ + new Uint32Array(1) + ]) as unknown as GPUVector<'float32'>; + const partitioned = createVector(fixture, 'partitioned-residual', 'float32', [ + new Float32Array(1), + new Float32Array(0) + ]); + + expect(() => new LuGraphPageRank({...props, residual: wrongFormat})).toThrow( + /residual|float32|packed/ + ); + expect(() => new LuGraphPageRank({...props, residual: partitioned})).toThrow( + /residual|one|single|chunk/ + ); + }); + + test('accepts float32 score and residual slices at non-256-byte-aligned offsets', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture); + const output = createVector( + fixture, + 'offset-page-rank', + 'float32', + [new Float32Array(props.topology.graph.vertexCount)], + {byteOffset: 4} + ); + const residual = createVector(fixture, 'offset-residual', 'float32', [new Float32Array(1)], { + byteOffset: 4 + }); + + const pageRank = new LuGraphPageRank({...props, output, residual}); + expect(pageRank.output.data[0].byteOffset).toBe(4); + expect(pageRank.residual?.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 score output backed by an existing physical allocation: %s', vectorName => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {vertexCount: 1, weighted: true, residual: true}); + const vector = getTopologyVector(props.topology, vectorName); + const output = createVector(fixture, 'aliased-page-rank', 'float32', [new Float32Array(1)], { + buffer: vector.data[0].buffer + }); + + expect(() => new LuGraphPageRank({...props, output})).toThrow( + /output|distinct|physical|allocation/ + ); + }); + + test('rejects residual status aliasing scores or topology status buffers', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(fixture, {vertexCount: 1, residual: true}); + const scoreAlias = createVector( + fixture, + 'aliased-score-residual', + 'float32', + [new Float32Array(1)], + {buffer: props.output.data[0].buffer} + ); + const topologyAlias = createVector( + fixture, + 'aliased-status-residual', + 'float32', + [new Float32Array(1)], + {buffer: props.topology.invalidEdgeCount.data[0].buffer} + ); + + expect(() => new LuGraphPageRank({...props, residual: scoreAlias})).toThrow( + /residual|distinct|physical|allocation/ + ); + expect(() => new LuGraphPageRank({...props, residual: topologyAlias})).toThrow( + /residual|distinct|physical|allocation/ + ); + }); + + test('unwraps borrowed DynamicBuffer views before checking physical score aliases', () => { + const fixture = createPageRankFixture(); + const props = createPageRankProps(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-page-rank', + 'float32', + [new Float32Array(props.topology.graph.vertexCount)], + {buffer: dynamicBuffer} + ); + + expect(() => new LuGraphPageRank({...props, output})).toThrow(/distinct|physical|allocation/); + expect(concreteBuffer.destroyed).toBe(false); + }); +}); + +function createPageRankFixture(): PageRankFixture { + const fixture = {device: new NullDevice({}), buffers: [], dynamicBuffers: [], vectors: []}; + pageRankFixtures.push(fixture); + return fixture; +} + +function createPageRankProps( + fixture: PageRankFixture, + options: { + vertexCount?: number; + directed?: boolean; + reverse?: boolean; + weighted?: boolean; + residual?: boolean; + } = {} +): LuGraphPageRankProps { + 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 directed = options.directed ?? true; + const graph = new LuGraph({ + vertexCount, + sourceVertices, + targetVertices, + edgeWeights, + edgeIds, + directed + }); + const forward = createAdjacency(fixture, 'forward', vertexCount, 5, options.weighted); + const includeReverse = options.reverse ?? directed; + const reverse = includeReverse + ? 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, 'pageRankScores', 'float32', [ + new Float32Array(vertexCount) + ]); + const residual = options.residual + ? createVector(fixture, 'pageRankResidual', 'float32', [new Float32Array(1)]) + : undefined; + return {topology, output, residual}; +} + +function createAdjacency( + fixture: PageRankFixture, + 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: PageRankFixture, + 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-page-rank.spec.ts b/modules/experimental/test/lugraph/lu-graph-page-rank.spec.ts new file mode 100644 index 0000000000..841ed81c07 --- /dev/null +++ b/modules/experimental/test/lugraph/lu-graph-page-rank.spec.ts @@ -0,0 +1,739 @@ +// 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, + LuGraphPageRank, + 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 { + addLuGraphPageRankToGraphWithDispatchLimit, + getLuGraphPageRankDispatchLayout +} from '../../src/lugraph/lu-graph-page-rank-internals'; + +const SCORE_TOLERANCE = 2e-5; +const NORMALIZATION_TOLERANCE = 5e-5; +const RESIDUAL_TOLERANCE = 8e-5; + +type ScalarFormat = 'uint32' | 'float32'; + +type PageRankScenario = { + name: string; + vertexCount: number; + sourceChunks: number[][]; + targetChunks: number[][]; + weightChunks?: number[][]; + directed?: boolean; + damping?: number; + iterations?: number; + residual?: boolean; + capacity?: number; + reverseCapacity?: number; + maximumWorkgroups?: number; + byteOffset?: number; +}; + +type ExpectedPageRank = { + scores: number[]; + residual: number; + invalidEdgeCount: number; + forwardCount: number; + reverseCount: number; + forwardOverflow: boolean; + reverseOverflow: boolean; + failed: boolean; +}; + +type PageRankExecutionFixture = { + device: Device; + buffers: Buffer[]; + vectors: GPUVector[]; + graph: LuGraph; + topology: LuGraphTopology; + pageRank: LuGraphPageRank; + commandGraph: GPUCommandGraph; + compiled?: ReturnType; +}; + +const pageRankScenarios: PageRankScenario[] = [ + { + name: 'empty directed graphs publish zero final residual and no rank rows', + vertexCount: 0, + sourceChunks: [], + targetChunks: [], + capacity: 0, + reverseCapacity: 0, + iterations: 2 + }, + { + name: 'empty undirected graphs support omitted residual output', + vertexCount: 0, + sourceChunks: [], + targetChunks: [], + capacity: 0, + directed: false, + residual: false, + iterations: 2 + }, + { + name: 'one isolated dangling vertex retains all normalized rank mass', + vertexCount: 1, + sourceChunks: [[]], + targetChunks: [[]], + capacity: 0, + reverseCapacity: 0, + iterations: 3 + }, + { + name: 'all dangling vertices uniformly redistribute probability without residual', + vertexCount: 5, + sourceChunks: [[], []], + targetChunks: [[], []], + capacity: 0, + reverseCapacity: 0, + iterations: 5 + }, + { + name: 'a directed dangling chain matches normalized reverse-pull PageRank', + vertexCount: 4, + sourceChunks: [[0, 1], [], [2]], + targetChunks: [[1, 2], [], [3]], + damping: 0.85, + iterations: 8 + }, + { + name: 'directed cycles retain symmetric uniform stationary rank', + vertexCount: 4, + sourceChunks: [[0, 1], [], [2, 3]], + targetChunks: [[1, 2], [], [3, 0]], + iterations: 6 + }, + { + name: 'incoming-star importance confirms reverse adjacency pull direction', + vertexCount: 5, + sourceChunks: [[1, 2], [], [3, 4]], + targetChunks: [[0, 0], [], [0, 0]], + damping: 0.9, + iterations: 10 + }, + { + name: 'disconnected groups and isolated nodes preserve normalized global probability', + vertexCount: 7, + sourceChunks: [[0, 1], [], [3, 4]], + targetChunks: [[1, 0], [], [4, 5]], + iterations: 7 + }, + { + name: 'duplicate edges and self-loops contribute through their exact outgoing degree', + vertexCount: 4, + sourceChunks: [[0, 0, 0], [], [1, 2]], + targetChunks: [[1, 1, 2], [], [1, 0]], + iterations: 8 + }, + { + name: 'zero damping immediately produces uniform ranks and zero final residual', + vertexCount: 5, + sourceChunks: [[0, 1, 3]], + targetChunks: [[1, 2, 4]], + damping: 0, + iterations: 3 + }, + { + name: 'unit damping redistributes dangling mass without division by zero', + vertexCount: 4, + sourceChunks: [[0, 1]], + targetChunks: [[1, 2]], + damping: 1, + iterations: 6 + }, + { + name: 'one iteration reports the exact final normalized L1 delta from uniform scores', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 1, 1]], + damping: 0.85, + iterations: 1 + }, + { + name: 'optional final residual can be omitted without changing normalized scores', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + residual: false, + iterations: 5 + }, + { + name: 'float32 edge attributes are preserved but links are intentionally unweighted', + vertexCount: 5, + sourceChunks: [[0, 0], [], [2, 3]], + targetChunks: [[1, 2], [], [1, 4]], + weightChunks: [[0.5, 20], [], [4, 8]], + iterations: 7 + }, + { + name: 'undirected graphs reuse symmetric forward CSR without reverse adjacency', + vertexCount: 5, + sourceChunks: [[0, 0], [], [2, 3]], + targetChunks: [[1, 2], [], [3, 4]], + directed: false, + iterations: 8 + }, + { + name: 'invalid endpoints become excluded links and newly dangling vertices', + vertexCount: 5, + sourceChunks: [[0, 9], [], [2, 3, 4]], + targetChunks: [[1, 2], [], [8, 4, 4]], + iterations: 7 + }, + { + name: 'forward CSR overflow fails closed with zero scores and zero residual', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + capacity: 0, + iterations: 3 + }, + { + name: 'reverse CSR overflow also fails closed with zero scores and residual', + vertexCount: 4, + sourceChunks: [[0, 1, 2]], + targetChunks: [[1, 2, 3]], + reverseCapacity: 1, + iterations: 3 + }, + { + name: 'undirected forward overflow fails closed without requiring a reverse status', + vertexCount: 4, + sourceChunks: [[0, 1]], + targetChunks: [[1, 2]], + directed: false, + capacity: 2, + iterations: 3 + }, + { + name: 'non-256-aligned CSR, score, and residual views preserve float32 binding offsets', + vertexCount: 5, + sourceChunks: [[0, 1, 3]], + targetChunks: [[1, 2, 4]], + iterations: 4, + byteOffset: 4 + }, + { + name: 'bounded 3D pull and hierarchical dangling reductions process 1025 vertices', + vertexCount: 1025, + sourceChunks: [ + Array.from({length: 600}, (_, vertexIndex) => vertexIndex), + [], + Array.from({length: 424}, (_, vertexIndex) => vertexIndex + 600) + ], + targetChunks: [ + Array.from({length: 600}, (_, vertexIndex) => vertexIndex + 1), + [], + Array.from({length: 424}, (_, vertexIndex) => vertexIndex + 601) + ], + iterations: 2, + maximumWorkgroups: 2 + } +]; + +test('LuGraphPageRank plans bounded three-dimensional pull and reduction dispatch', tapeTest => { + tapeTest.deepEqual(getLuGraphPageRankDispatchLayout(0, 2), {x: 1, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphPageRankDispatchLayout(512, 2), {x: 2, y: 1, z: 1}); + tapeTest.deepEqual(getLuGraphPageRankDispatchLayout(513, 2), {x: 2, y: 2, z: 1}); + tapeTest.deepEqual(getLuGraphPageRankDispatchLayout(1025, 2), {x: 2, y: 2, z: 2}); + tapeTest.throws(() => getLuGraphPageRankDispatchLayout(2049, 2), /3D dispatch limit/); + tapeTest.end(); +}); + +for (const scenario of pageRankScenarios) { + test(`LuGraphPageRank GPU analytics: ${scenario.name}`, async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const expected = calculateExpectedPageRank(scenario); + const fixture = createExecutionFixture(device, scenario, expected); + try { + compilePageRank(fixture, scenario.maximumWorkgroups); + executePageRank(fixture); + await assertPageRank(tapeTest, fixture, expected); + tapeTest.deepEqual( + fixture.graph.sourceVertices.data.map(chunk => chunk.length), + scenario.sourceChunks.map(chunk => chunk.length), + 'rank evaluation preserves caller-owned source chunks and empty edge batches' + ); + } finally { + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); + }); +} + +test('LuGraphPageRank reinitializes normalized scores after source updates without hidden GPU work', async tapeTest => { + const device = await getWebGPUTestDevice(); + if (!device) { + tapeTest.comment('WebGPU is not available'); + tapeTest.end(); + return; + } + + const original: PageRankScenario = { + name: 'repeated PageRank encoding', + vertexCount: 6, + sourceChunks: [[0, 1], [], [2, 4]], + targetChunks: [[1, 2], [], [3, 5]], + damping: 0.85, + iterations: 6 + }; + const fixture = createExecutionFixture(device, original, calculateExpectedPageRank(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 { + compilePageRank(fixture); + tapeTest.equal( + submitSpy.mock.calls.length, + 0, + 'topology and rank construction never submit work' + ); + tapeTest.ok( + sourceReadbackSpies.every(spy => spy.mock.calls.length === 0), + 'PageRank never reads source edge columns back to the CPU' + ); + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + + executePageRank(fixture); + await assertPageRank(tapeTest, fixture, calculateExpectedPageRank(original)); + + const sourceBuffer = fixture.graph.sourceVertices.data[0].buffer as Buffer; + sourceBuffer.write(Uint32Array.from([9, 1])); + const updated = {...original, sourceChunks: [[9, 1], [], [2, 4]]}; + executePageRank(fixture); + await assertPageRank(tapeTest, fixture, calculateExpectedPageRank(updated)); + tapeTest.equal( + fixture.graph.sourceVertices.data[0].buffer, + sourceBuffer, + 'repeated ranking preserves exact source chunk and physical buffer identity' + ); + } finally { + submitSpy.mockRestore(); + for (const sourceReadbackSpy of sourceReadbackSpies) sourceReadbackSpy.mockRestore(); + destroyExecutionFixture(tapeTest, fixture); + } + + tapeTest.end(); +}); + +/** Computes unweighted PageRank with exact dangling redistribution and per-step normalization. */ +function calculateExpectedPageRank(scenario: PageRankScenario): ExpectedPageRank { + const outgoing = Array.from({length: scenario.vertexCount}, () => [] as number[]); + const incoming = Array.from({length: scenario.vertexCount}, () => [] as number[]); + let invalidEdgeCount = 0; + let validEdgeCount = 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++; + outgoing[source].push(target); + incoming[target].push(source); + if (scenario.directed === false && source !== target) { + outgoing[target].push(source); + incoming[source].push(target); + } + } + } + + const forwardCount = outgoing.reduce((count, neighbors) => count + neighbors.length, 0); + const reverseCount = scenario.directed === false ? 0 : validEdgeCount; + const forwardOverflow = forwardCount > (scenario.capacity ?? forwardCount); + const reverseOverflow = + scenario.directed !== false && reverseCount > (scenario.reverseCapacity ?? reverseCount); + const failed = forwardOverflow || reverseOverflow; + let scores = new Array(scenario.vertexCount).fill(0); + let residual = 0; + + if (!failed && scenario.vertexCount > 0) { + scores.fill(1 / scenario.vertexCount); + const damping = scenario.damping ?? 0.85; + const iterations = scenario.iterations ?? 40; + for (let iteration = 0; iteration < iterations; iteration++) { + const danglingMass = scores.reduce( + (sum, score, vertexIndex) => sum + (outgoing[vertexIndex].length === 0 ? score : 0), + 0 + ); + const next = incoming.map(neighbors => { + const contribution = neighbors.reduce( + (sum, neighbor) => sum + scores[neighbor] / outgoing[neighbor].length, + 0 + ); + return ( + (1 - damping) / scenario.vertexCount + + damping * (contribution + danglingMass / scenario.vertexCount) + ); + }); + const mass = next.reduce((sum, score) => sum + score, 0); + const normalized = next.map(score => (mass > 0 ? score / mass : 0)); + residual = normalized.reduce( + (difference, score, vertexIndex) => difference + Math.abs(score - scores[vertexIndex]), + 0 + ); + scores = normalized; + } + } + + return { + scores, + residual, + invalidEdgeCount, + forwardCount, + reverseCount, + forwardOverflow, + reverseOverflow, + failed + }; +} + +function createExecutionFixture( + device: Device, + scenario: PageRankScenario, + expected: ExpectedPageRank +): PageRankExecutionFixture { + 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 directed = scenario.directed ?? true; + const graph = new LuGraph({ + vertexCount: scenario.vertexCount, + sourceVertices, + targetVertices, + edgeWeights, + directed + }); + const forward = createOutputAdjacency( + device, + buffers, + vectors, + 'forward', + scenario.vertexCount, + scenario.capacity ?? expected.forwardCount, + Boolean(edgeWeights), + scenario.byteOffset + ); + const reverse = directed + ? 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, + 'page-rank-scores', + 'float32', + scenario.vertexCount, + scenario.byteOffset + ); + const residual = + scenario.residual === false + ? undefined + : createOutputVector( + device, + buffers, + vectors, + 'page-rank-residual', + 'float32', + 1, + scenario.byteOffset + ); + const pageRank = new LuGraphPageRank({ + topology, + output, + damping: scenario.damping, + iterations: scenario.iterations, + residual + }); + + return { + device, + buffers, + vectors, + graph, + topology, + pageRank, + 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 compilePageRank(fixture: PageRankExecutionFixture, maximumWorkgroups?: number): void { + fixture.topology.addToGraph(fixture.commandGraph); + if (maximumWorkgroups === undefined) { + fixture.pageRank.addToGraph(fixture.commandGraph); + } else { + addLuGraphPageRankToGraphWithDispatchLimit( + fixture.pageRank, + fixture.commandGraph, + maximumWorkgroups + ); + } + fixture.compiled = fixture.commandGraph.compile(); +} + +function executePageRank(fixture: PageRankExecutionFixture): void { + const commandEncoder = fixture.device.createCommandEncoder({id: 'lu-graph-page-rank-test'}); + fixture.compiled!.encode(commandEncoder, {parameters: undefined}); + fixture.device.submit(commandEncoder.finish()); +} + +async function assertPageRank( + tapeTest: Test, + fixture: PageRankExecutionFixture, + expected: ExpectedPageRank +): Promise { + const [scores, residual, invalidEdgeCount, forwardOverflow, reverseOverflow] = await Promise.all([ + readFloat32Vector(fixture.pageRank.output), + fixture.pageRank.residual + ? readFloat32Vector(fixture.pageRank.residual) + : Promise.resolve(undefined), + readUint32Vector(fixture.topology.invalidEdgeCount), + readUint32Vector(fixture.topology.forward.overflow), + fixture.topology.reverse + ? readUint32Vector(fixture.topology.reverse.overflow) + : Promise.resolve(undefined) + ]); + + tapeTest.equal( + scores.length, + expected.scores.length, + 'one float32 score is published per vertex' + ); + const largestScoreError = scores.reduce( + (largest, score, vertexIndex) => + Math.max(largest, Math.abs(score - expected.scores[vertexIndex])), + 0 + ); + tapeTest.ok( + largestScoreError <= SCORE_TOLERANCE, + `float32 reverse-pull ranks match the CPU oracle within ${SCORE_TOLERANCE}` + ); + if (!expected.failed && scores.length > 0) { + tapeTest.ok( + scores.every(score => Number.isFinite(score) && score >= 0), + 'every score is finite and nonnegative' + ); + const rankMass = scores.reduce((sum, score) => sum + score, 0); + tapeTest.ok( + Math.abs(rankMass - 1) <= NORMALIZATION_TOLERANCE, + 'dangling redistribution and float32 normalization preserve unit rank mass' + ); + } + if (residual) { + tapeTest.ok( + Math.abs(residual[0] - expected.residual) <= RESIDUAL_TOLERANCE, + 'optional residual equals the final normalized L1 PageRank delta' + ); + } + tapeTest.equal( + invalidEdgeCount[0], + expected.invalidEdgeCount, + 'invalid graph edges are excluded' + ); + tapeTest.equal( + forwardOverflow[0], + Number(expected.forwardOverflow), + 'forward capacity remains explicit' + ); + if (reverseOverflow) { + tapeTest.equal( + reverseOverflow[0], + Number(expected.reverseOverflow), + 'reverse capacity remains explicit' + ); + } + if (expected.failed) { + tapeTest.ok( + scores.every(score => score === 0), + 'required CSR overflow fails closed with zero scores' + ); + if (residual) tapeTest.equal(residual[0], 0, 'failed topology publishes a zero final residual'); + } +} + +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)); +} + +async function readFloat32Vector(vector: GPUVector<'float32'>): Promise { + if (vector.length === 0) return []; + const data = vector.data[0]; + const bytes = await (data.buffer as Buffer).readAsync( + data.byteOffset, + vector.length * Float32Array.BYTES_PER_ELEMENT + ); + return Array.from(new Float32Array(bytes.buffer, bytes.byteOffset, vector.length)); +} + +function destroyExecutionFixture(tapeTest: Test, fixture: PageRankExecutionFixture): void { + fixture.compiled?.destroy(); + for (const vector of fixture.vectors) vector.destroy(); + tapeTest.ok( + fixture.buffers.every(buffer => !buffer.destroyed), + 'graph-owned PageRank reduction scratch never destroys caller-owned physical buffers' + ); + for (const buffer of fixture.buffers) buffer.destroy(); +} diff --git a/test/examples/lugraph-docs.node.spec.ts b/test/examples/lugraph-docs.node.spec.ts new file mode 100644 index 0000000000..8f220526bc --- /dev/null +++ b/test/examples/lugraph-docs.node.spec.ts @@ -0,0 +1,108 @@ +// luma.gl +// SPDX-License-Identifier: MIT +// SPDX-FileCopyrightText: Copyright (c) vis.gl contributors + +import {readFileSync} from 'node:fs'; +import {describe, expect, test} from 'vitest'; + +const graphDocumentation = readFileSync( + new URL('../../docs/api-reference/experimental/lugraph.md', import.meta.url), + 'utf8' +); +const packageDocumentation = readFileSync( + new URL('../../modules/experimental/src/lugraph/README.md', import.meta.url), + 'utf8' +); +const experimentalOverview = readFileSync( + new URL('../../docs/api-reference/experimental/README.md', import.meta.url), + 'utf8' +); +const sidebar = readFileSync(new URL('../../docs/table-of-contents.json', import.meta.url), 'utf8'); +const experimentalTabs = readFileSync( + new URL('../../website/src/components/docs/experimental-docs-tabs.tsx', import.meta.url), + 'utf8' +); + +describe('luGraph GPU-resident graph analytics documentation', () => { + test('publishes one canonical guide in both experimental sidebars, overview, and tabs', () => { + expect(graphDocumentation).toContain('# luGraph: GPU-Resident Graph Analytics'); + expect(graphDocumentation).toContain(''); + expect(sidebar.match(/"api-reference\/experimental\/lugraph"/gu)).toHaveLength(2); + expect(experimentalTabs).toContain("| 'lugraph'"); + expect(experimentalTabs).toContain("href: '/docs/api-reference/experimental/lugraph'"); + expect(experimentalOverview).toContain('## GPU-resident Graph Analytics'); + expect(experimentalOverview).toContain('/docs/api-reference/experimental/lugraph'); + expect(packageDocumentation).toContain('/docs/api-reference/experimental/lugraph'); + }); + + test('explains graph motivation, appropriate workloads, and concrete application use cases', () => { + expect(graphDocumentation).toContain('## Overview'); + expect(graphDocumentation).toContain('## Why keep a graph on the GPU?'); + expect(graphDocumentation).toContain('## When should I use luGraph?'); + expect(graphDocumentation).toContain('Social and communication networks'); + expect(graphDocumentation).toContain('Software and service dependencies'); + expect(graphDocumentation).toContain('Transaction and fraud investigations'); + expect(graphDocumentation).toContain('Transport and infrastructure maps'); + expect(graphDocumentation).toContain('Knowledge and citation graphs'); + expect(graphDocumentation).toContain('A small, CPU-resident, one-off analysis'); + expect(graphDocumentation).toContain('**Question: How many direct relationships'); + expect(graphDocumentation).toContain('**Question: Which entities can I reach'); + expect(graphDocumentation).toContain('**Question: Which vertices belong to the same connected'); + expect(graphDocumentation).toContain('**Question: Which vertices receive influence'); + }); + + test('introduces every available operation and composes its actual optional entry point', () => { + for (const graphOperation of [ + 'LuGraph', + 'LuGraphTopology', + 'LuGraphDegree', + 'LuGraphBreadthFirstSearch', + 'LuGraphConnectedComponents', + 'LuGraphPageRank' + ]) { + expect(graphDocumentation, graphOperation).toContain(graphOperation); + } + + expect(packageDocumentation).toContain('compressed adjacency'); + expect(packageDocumentation).toContain('vertex-degree queries'); + expect(packageDocumentation).toContain('breadth-first shortest paths'); + expect(packageDocumentation).toContain('weakly connected components'); + expect(packageDocumentation).toContain('normalized PageRank'); + expect(graphDocumentation).toContain("from '@luma.gl/experimental/lugraph';"); + expect(graphDocumentation).toContain('topology.addToGraph(workflow);'); + expect(graphDocumentation).toContain('const compiled = workflow.compile();'); + expect(graphDocumentation).toContain('compiled.encode(encoder, {parameters: undefined});'); + expect(graphDocumentation).toContain('device.submit(encoder.finish());'); + }); + + test('documents overflow, direction, probability, iteration, and ownership boundaries honestly', () => { + expect(graphDocumentation).toContain('`vertexCount + 1` rows'); + expect(graphDocumentation).toContain( + 'Neighbor order within each vertex is intentionally unspecified' + ); + expect(graphDocumentation).toContain('Degrees come from complete CSR offsets'); + expect(graphDocumentation).toContain('Directed weak components use forward adjacency'); + expect(graphDocumentation).toContain('Directed graphs require reverse CSR'); + expect(graphDocumentation).toContain('dangling vertices with no outgoing edges'); + expect(graphDocumentation).toContain('default damping is `0.85`'); + expect(graphDocumentation).toContain('85% chance of following an outgoing link'); + expect(graphDocumentation).toContain('15% chance of jumping to a uniformly chosen vertex'); + expect(graphDocumentation).toContain('default bounded iteration count is `40`'); + expect(graphDocumentation).toContain("final iteration's L1 score change"); + expect(graphDocumentation).toContain('absolute differences between the last two normalized'); + expect(graphDocumentation).toContain('not an automatic convergence threshold'); + expect(graphDocumentation).toContain('physically distinct GPU buffer allocations'); + expect(graphDocumentation).toContain('does not imply distributed or multi-GPU execution'); + }); + + test('preserves independent MIT ownership and accurate NVIDIA RAPIDS inspiration', () => { + for (const documentation of [graphDocumentation, packageDocumentation]) { + expect(documentation).toContain('NVIDIA RAPIDS cuGraph'); + expect(documentation).toContain('https://github.com/rapidsai/cugraph'); + expect(documentation).toContain('Apache License 2.0'); + expect(documentation).toContain('MIT-licensed'); + expect(documentation).toContain('does not copy or translate cuGraph source code'); + expect(documentation).toMatch(/endorse(?:d|ment)/u); + } + }); +}); diff --git a/website/src/components/docs/experimental-docs-tabs.tsx b/website/src/components/docs/experimental-docs-tabs.tsx index 232a7481a3..074e857826 100644 --- a/website/src/components/docs/experimental-docs-tabs.tsx +++ b/website/src/components/docs/experimental-docs-tabs.tsx @@ -10,6 +10,7 @@ export type ExperimentalDocsTabId = | 'deferred-scene-renderer' | 'pbr-environment' | 'luproj' + | 'lugraph' | 'luxfilter' | 'lutrace' | 'g-buffer' @@ -42,6 +43,7 @@ const EXPERIMENTAL_DOCS_TABS: ExperimentalDocsTab[] = [ href: '/docs/api-reference/experimental/pbr-environment' }, {id: 'luproj', label: 'GPU Projection', href: '/docs/api-reference/experimental/luproj'}, + {id: 'lugraph', label: 'GPU Graphs', href: '/docs/api-reference/experimental/lugraph'}, {id: 'luxfilter', label: 'LuxFilter', href: '/docs/api-reference/experimental/luxfilter'}, {id: 'lutrace', label: 'GPU Traces', href: '/docs/api-reference/experimental/lutrace'}, {id: 'g-buffer', label: 'GBuffer', href: '/docs/api-reference/experimental/g-buffer'},