2a:null 2b:["$","title",null,{"children":"SSSP | AntV | Observable"}] 30:T7b5,## Kernel Code ```cuda __global__ void relax(int N, int MAX_VAL, int *d_in_V, int *d_in_I, int *d_in_E, int *d_in_W, int *d_out_D, int *d_out_Di) { unsigned int index = threadIdx.x + (blockDim.x * blockIdx.x); if (index < N - 1) { // do index < N - 1 because nth element of I array points to the end of E array for (int j = d_in_I[index]; j < d_in_I[index + 1]; j++) { int w = d_in_W[j]; int du = d_out_D[index]; int dv = d_out_D[d_in_E[j]]; int newDist = du + w; // Check if the distance is already set to max then just take the max since, // Cuda implementation gives this when a number is added to already max value of int. // E.g 2147483647 + 5 becomes -2147483644 if (du == MAX_VAL){ newDist = MAX_VAL; } if (newDist < dv) { atomicMin(&d_out_Di[d_in_E[j]],newDist); } } } } ``` ```wgsl struct Buffer { data: array, }; struct AtomicBuffer { data: array>, }; @group(0) @binding(0) var d_in_E : Buffer; @group(0) @binding(1) var d_in_I : Buffer; @group(0) @binding(2) var d_in_W : Buffer; @group(0) @binding(3) var d_out_D : Buffer; @group(0) @binding(4) var d_out_Di : AtomicBuffer; @compute @workgroup_size(\${BLOCKS}, \${BLOCK_SIZE}) fn main( @builtin(global_invocation_id) global_id : vec3 ) { var index = global_id.x; if (index < \${V.length}u) { for (var j = d_in_I.data[index]; j < d_in_I.data[index + 1u]; j = j + 1) { var w = d_in_W.data[j]; var du = d_out_D.data[index]; var dv = d_out_D.data[d_in_E.data[j]]; var newDist = du + w; if (du == \${MAX_DISTANCE}) { newDist = \${MAX_DISTANCE}; } if (newDist < dv) { atomicMin(&d_out_Di.data[d_in_E.data[j]], newDist); } } } } ```31:T1dc4,result = (async () => { const $canvas = createCanvas(); const deviceContribution = createDeviceContributionWebGPU(); const swapChain = await deviceContribution.createSwapChain($canvas); swapChain.configureSwapChain($canvas.width, $canvas.height); const device = swapChain.getDevice(); const { BufferUsage, BufferFrequencyHint, VertexStepMode, Format, PrimitiveTopology, TextureUsage, TransparentWhite, TransparentBlack, ChannelWriteMask, BlendMode, BlendFactor, CullMode, CompareMode, WrapMode, TexFilterMode, MipFilterMode } = deviceAPI; const sourceId = "A"; const weightPropertyName = ""; const maxDistance = 1000000; // The total number of workgroup invocations (4096) exceeds the maximum allowed (256). const BLOCK_SIZE = 1; const BLOCKS = 256; const MAX_DISTANCE = maxDistance; const graph = new Graph.Graph(simpleDataset); const { V, E, I, nodeId2IndexMap, edges } = convertGraphData2CSC(graph); let W; const sourceIdx = nodeId2IndexMap[sourceId]; if (weightPropertyName) { W = edges.map((edgeConfig) => Number(edgeConfig.data[weightPropertyName])); } else { // all the vertex has the same weight W = new Array(E.length).fill(1); } const VBuffer = device.createBuffer({ usage: BufferUsage.STORAGE, viewOrSize: new Int32Array(V) }); const EBuffer = device.createBuffer({ usage: BufferUsage.STORAGE, viewOrSize: new Int32Array(E) }); const IBuffer = device.createBuffer({ usage: BufferUsage.STORAGE, viewOrSize: new Int32Array(I) }); const WBuffer = device.createBuffer({ usage: BufferUsage.STORAGE, viewOrSize: new Int32Array(W) }); // mark source vertex const view = new Array(V.length).fill(MAX_DISTANCE); view[sourceIdx] = 0; const DOutBuffer = device.createBuffer({ usage: BufferUsage.STORAGE | BufferUsage.COPY_SRC, viewOrSize: new Int32Array(view) }); const DiOutBuffer = device.createBuffer({ usage: BufferUsage.STORAGE | BufferUsage.COPY_SRC, viewOrSize: new Int32Array(view) }); // store predecessors const POutBuffer = device.createBuffer({ usage: BufferUsage.STORAGE | BufferUsage.COPY_SRC, viewOrSize: new Int32Array(view) }); const readback = device.createReadback(); const relaxProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; struct AtomicBuffer { data: array>, }; @group(0) @binding(0) var d_in_E : Buffer; @group(0) @binding(1) var d_in_I : Buffer; @group(0) @binding(2) var d_in_W : Buffer; @group(0) @binding(3) var d_out_D : Buffer; @group(0) @binding(4) var d_out_Di : AtomicBuffer; @compute @workgroup_size(${BLOCKS}, ${BLOCK_SIZE}) fn main( @builtin(global_invocation_id) global_id : vec3 ) { var index = global_id.x; if (index < ${V.length}u) { for (var j = d_in_I.data[index]; j < d_in_I.data[index + 1u]; j = j + 1) { var w = d_in_W.data[j]; var du = d_out_D.data[index]; var dv = d_out_D.data[d_in_E.data[j]]; var newDist = du + w; if (du == ${MAX_DISTANCE}) { newDist = ${MAX_DISTANCE}; } if (newDist < dv) { atomicMin(&d_out_Di.data[d_in_E.data[j]], newDist); } } } }` } }); const updateDistanceProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; @group(0) @binding(0) var d_out_D : Buffer; @group(0) @binding(1) var d_out_Di : Buffer; @compute @workgroup_size(${BLOCKS}, ${BLOCK_SIZE}) fn main( @builtin(global_invocation_id) global_id : vec3 ) { var index = global_id.x; if (index < ${V.length}u) { if (d_out_D.data[index] > d_out_Di.data[index]) { d_out_D.data[index] = d_out_Di.data[index]; } d_out_Di.data[index] = d_out_D.data[index]; } } ` } }); const updatePredProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; struct AtomicBuffer { data: array>, }; @group(0) @binding(0) var d_in_V : Buffer; @group(0) @binding(1) var d_in_E : Buffer; @group(0) @binding(2) var d_in_I : Buffer; @group(0) @binding(3) var d_in_W : Buffer; @group(0) @binding(4) var d_out_D : Buffer; @group(0) @binding(5) var d_out_P : AtomicBuffer; @compute @workgroup_size(${BLOCKS}, ${BLOCK_SIZE}) fn main( @builtin(global_invocation_id) global_id : vec3 ) { var index = global_id.x; if (index < ${V.length}u) { for (var j = d_in_I.data[index]; j < d_in_I.data[index + 1u]; j = j + 1) { var u = d_in_V.data[index]; var w = d_in_W.data[j]; var dis_u = d_out_D.data[index]; var dis_v = d_out_D.data[d_in_E.data[j]]; if (dis_v == dis_u + w) { atomicMin(&d_out_P.data[d_in_E.data[j]], u); } } } } ` } }); const relaxPipeline = device.createComputePipeline({ inputLayout: null, program: relaxProgram }); const updateDistancePipeline = device.createComputePipeline({ inputLayout: null, program: updateDistanceProgram }); const updatePredPipeline = device.createComputePipeline({ inputLayout: null, program: updatePredProgram }); const relaxBindings = device.createBindings({ pipeline: relaxPipeline, storageBufferBindings: [ { buffer: EBuffer }, { buffer: IBuffer }, { buffer: WBuffer }, { buffer: DOutBuffer }, { buffer: DiOutBuffer } ] }); const updateDistanceBindings = device.createBindings({ pipeline: updateDistancePipeline, storageBufferBindings: [ { buffer: DOutBuffer }, { buffer: DiOutBuffer } ] }); const updatePredBindings = device.createBindings({ pipeline: updatePredPipeline, storageBufferBindings: [ { buffer: VBuffer }, { buffer: EBuffer }, { buffer: IBuffer }, { buffer: WBuffer }, { buffer: DOutBuffer }, { buffer: POutBuffer } ] }); const computePass = device.createComputePass(); const grids = Math.ceil(V.length / (BLOCKS * BLOCK_SIZE)); for (let i = 1; i < V.length; i++) { computePass.setPipeline(relaxPipeline); computePass.setBindings(relaxBindings); computePass.dispatchWorkgroups(grids); computePass.setPipeline(updateDistancePipeline); computePass.setBindings(updateDistanceBindings); computePass.dispatchWorkgroups(grids); } computePass.setPipeline(updatePredPipeline); computePass.setBindings(updatePredBindings); computePass.dispatchWorkgroups(grids); device.submitPass(computePass); const out = await readback.readBuffer(DiOutBuffer, 0, new Int32Array(view)); const predecessor = await readback.readBuffer( POutBuffer, 0, new Int32Array(view) ); invalidation.then(() => { // storeProgram.destroy(); // storePipeline.destroy(); // matmulProgram.destroy(); // matmulPipeline.destroy(); // rankDiffProgram.destroy(); // rankDiffPipeline.destroy(); device.destroy(); // For debug. device.checkForLeaks(); }); return Array.from(out).map((distance, i) => ({ target: graph.getAllNodes()[V[i]].id, distance, predecessor: graph.getAllNodes()[V[predecessor[i]]]?.id })); })()32:T6c7,simpleDataset = { return { nodes: [ { id: "A" }, { id: "B" }, { id: "C" }, { id: "D" }, { id: "E" }, { id: "F" }, { id: "G" }, { id: "H" }, { id: "I" }, { id: "J" }, { id: "K" } ], edges: [ { id: "e1", source: "D", target: "A" }, { id: "e2", source: "D", target: "B" }, { id: "e3", source: "B", target: "C" }, { id: "e4", source: "C", target: "B" }, { id: "e5", source: "F", target: "B" }, { id: "e6", source: "F", target: "E" }, { id: "e7", source: "E", target: "F" }, { id: "e8", source: "E", target: "D" }, { id: "e9", source: "E", target: "B" }, { id: "e10", source: "K", target: "E" }, { id: "e11", source: "J", target: "E" }, { id: "e12", source: "I", target: "E" }, { id: "e13", source: "H", target: "E" }, { id: "e14", source: "G", target: "E" }, { id: "e15", source: "G", target: "B" }, { id: "e16", source: "H", target: "B" }, { id: "e17", source: "I", target: "B" } ] }; }33:[] 2e:{"id":"e3dd82621880eb5d","slug":"sssp","title":"SSSP","thumbnail":"c48907b5e47627ccac73580e4145dae9c8192d742a0d82602de99c4e72ab6ce6","access_level":"public","create_time":"$D2023-12-23T11:57:14.486Z","update_time":"$D2023-12-23T12:18:00.303Z","view_time":null,"first_public_time":"$D2023-12-23T11:57:14.486Z","fork":{"id":"da8420861f120e83","slug":"page-rank","title":"Page Rank","owner":{"id":"96a4168c70a60e69","login":"antv","name":"AntV"}},"stars":0,"starred":false,"owner":{"type":"team","id":"96a4168c70a60e69","login":"antv","name":"AntV","avatar_url":"https://avatars-next.observableusercontent.com/9cfb1b3742d5fee0ae1105c158818aad1b91fd02751860be879c813cbaac3abf","tier":"starter_2024"},"version":252,"publish_version":null,"role":null,"authored":false,"authors":[{"type":"individual","id":"bf67d760eefcbffd","login":"xiaoiver","name":"","avatar_url":"https://avatars-next.observableusercontent.com/f4a568df4cb69f834e43d657ac3e7239ed96aac8fb20e8c31c9c1ddcdf8b54a0","tier":null}],"members":[{"type":"team","id":"96a4168c70a60e69","login":"antv","name":"AntV","avatar_url":"https://avatars-next.observableusercontent.com/9cfb1b3742d5fee0ae1105c158818aad1b91fd02751860be879c813cbaac3abf","tier":"starter_2024","role":"editor"},{"type":"individual","id":"bf67d760eefcbffd","login":"xiaoiver","name":"","avatar_url":"https://avatars-next.observableusercontent.com/f4a568df4cb69f834e43d657ac3e7239ed96aac8fb20e8c31c9c1ddcdf8b54a0","tier":null,"role":"editor"}],"comments":[],"custom_thumbnail":false,"trashed":false,"subscribed":null,"forks":1,"collection":null,"type":"notebook","body":{"title":"SSSP","cells":[{"id":0,"value":"# SSSP\n\nMigrate the CUDA algorithm to WebGPU.\n\n* https://www.lewuathe.com/illustration-of-distributed-bellman-ford-algorithm.html\n* https://github.com/sengorajkumar/gpu_graph_algorithms\n* https://docs.rapids.ai/api/cugraph/stable/api_docs/api/cugraph.traversal.sssp.sssp.html\n\nThe Single Source Shortest Path (SSSP) algorithm, as the name suggests finds the shortest path from a source node to all other nodes. We use Bellman-ford implementation (compared to Dijkstra supporting negatively weighted edges) with complexity O(V * E).","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":195,"value":"## Store the Graph in CSC format\n\nThe first problem we have to solve is how to store the nodes and edges in the graph.\n\nDescribing a graph structure in JS is very flexible, we can use an object like `{ vertices: {}, edges: {} }`.\n\nBut the GPU only recognizes primitive memory structures such as `[1, 2, 3, 4]` (the same is true for WebGL-based solutions storing them as textures). nvGRAPH gives three ways to store a graph with such a linear structure, and different structures are suitable for different algorithms:\n\n```\ntypedef enum\n{\n NVGRAPH_CSR_32 = 0, SrSPMV\n NVGRAPH_CSC_32 = 1, SSSP, WidestPath, Pagerank\n NVGRAPH_COO_32 = 2\n} nvgraphTopologyType_t;\n```\n\nWe choose CSC (Compressed Sparse Column) and use `V` and `E` to store the nodes and edges, `I` for the edge index (e.g. if we want to know how many edges there are from `A`, we can look up `AB` `AC` in `[0, 2]` of `E`), and `W` to store the weights of the edges:\n\n![CSC](https://pic3.zhimg.com/80/v2-231cfa44064c1e965044af7d9af9be1a_1440w.webp)","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":204,"value":"## Host Code","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":214,"value":"$30","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":211,"value":"## Result","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":17,"value":"$31","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined"},{"id":5,"value":"## Appendix","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":55,"value":"import {\n deviceAPI,\n createDeviceContributionWebGPU,\n createCanvas,\n createBlitPipelineAndBindings,\n registerShaderModule,\n createProgram,\n alias,\n prelude,\n glMatrix\n} from \"@antv/compute-toys\"","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined"},{"id":174,"value":"Graph = require(\"@antv/graphlib\")","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined"},{"id":178,"value":"$32","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined"},{"id":184,"value":"convertGraphData2CSC = (graphData) => {\n const V = [];\n const E = [];\n const From = [];\n const To = [];\n const I = [];\n const nodeId2IndexMap = {};\n const edges = [];\n graphData.getAllNodes().forEach((node, i) => {\n nodeId2IndexMap[node.id] = i;\n V.push(i);\n });\n\n let lastSource = \"\";\n let counter = 0;\n // sort by source\n [...graphData.getAllEdges()]\n .sort((a, b) => nodeId2IndexMap[a.source] - nodeId2IndexMap[b.source])\n .forEach((edgeConfig) => {\n const { source, target } = edgeConfig;\n edges.push(edgeConfig);\n E.push(nodeId2IndexMap[target]);\n From.push(nodeId2IndexMap[source]);\n To.push(nodeId2IndexMap[target]);\n\n if (source !== lastSource) {\n I.push(counter);\n lastSource = source;\n }\n\n counter++;\n });\n\n I.push(E.length);\n\n return {\n V,\n E,\n I,\n From,\n To,\n nodeId2IndexMap,\n edges\n };\n}","pinned":false,"hidden":false,"mode":"ojs","author":"$undefined","summary":"$undefined"}],"files":"$W33","stdlib":"1","titling":"auto"}} 2f:"$2e:body"