2b:null 2c:["$","title",null,{"children":"Page Rank | AntV | Observable"}] 31:T15cb,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 tolerance = 1e-5; const alpha = 0.85; let maxIterations = 1000; const BLOCK_SIZE = 1; const BLOCKS = 256; const { V, From, To } = convertGraphData2CSC(new Graph.Graph(simpleDataset)); const n = V.length; const graph = new Float32Array(new Array(n * n).fill((1 - alpha) / n)); const r = new Float32Array(new Array(n).fill(1 / n)); From.forEach((from, i) => { graph[To[i] * n + from] += alpha * 1.0; }); for (let j = 0; j < n; j++) { let sum = 0.0; for (let i = 0; i < n; ++i) { sum += graph[i * n + j]; } for (let i = 0; i < n; ++i) { if (sum != 0.0) { graph[i * n + j] /= sum; } else { graph[i * n + j] = 1 / n; } } } const rBuffer = device.createBuffer({ usage: BufferUsage.STORAGE | BufferUsage.COPY_SRC, viewOrSize: new Float32Array(r) }); const rLastBuffer = device.createBuffer({ usage: BufferUsage.STORAGE | BufferUsage.COPY_SRC, viewOrSize: new Float32Array(n) }); const graphBuffer = device.createBuffer({ usage: BufferUsage.STORAGE, viewOrSize: new Float32Array(graph) }); const readback = device.createReadback(); const storeProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; @group(2) @binding(0) var r : Buffer; @group(2) @binding(1) var r_last : 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) { r_last.data[index] = r.data[index]; } } ` } }); const matmulProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; @group(2) @binding(0) var graph : Buffer; @group(2) @binding(1) var r : Buffer; @group(2) @binding(2) var r_last : 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) { var sum = 0.0; for (var i = 0u; i < ${V.length}u; i = i + 1u) { sum = sum + r_last.data[i] * graph.data[index * ${V.length}u + i]; } r.data[index] = sum; } } ` } }); const rankDiffProgram = device.createProgram({ compute: { wgsl: ` struct Buffer { data: array, }; @group(2) @binding(0) var r : Buffer; @group(2) @binding(1) var r_last : 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) { r_last.data[index] = abs(r_last.data[index] - r.data[index]); } } ` } }); const storePipeline = device.createComputePipeline({ inputLayout: null, program: storeProgram }); const matmulPipeline = device.createComputePipeline({ inputLayout: null, program: matmulProgram }); const rankDiffPipeline = device.createComputePipeline({ inputLayout: null, program: rankDiffProgram }); const storeBindings = device.createBindings({ pipeline: storePipeline, storageBufferBindings: [ { buffer: rBuffer }, { buffer: rLastBuffer } ] }); const matmulBindings = device.createBindings({ pipeline: matmulPipeline, storageBufferBindings: [ { buffer: graphBuffer }, { buffer: rBuffer }, { buffer: rLastBuffer } ] }); const rankDiffBindings = device.createBindings({ pipeline: rankDiffPipeline, storageBufferBindings: [ { buffer: rBuffer }, { buffer: rLastBuffer } ] }); const grids = Math.ceil(V.length / (BLOCKS * BLOCK_SIZE)); const computePass = device.createComputePass(); while (maxIterations--) { computePass.setPipeline(storePipeline); computePass.setBindings(storeBindings); computePass.dispatchWorkgroups(grids); computePass.setPipeline(matmulPipeline); computePass.setBindings(matmulBindings); computePass.dispatchWorkgroups(grids); computePass.setPipeline(rankDiffPipeline); computePass.setBindings(rankDiffBindings); computePass.dispatchWorkgroups(grids); const last = await readback.readBuffer(rLastBuffer, 0, new Float32Array(n)); const result = last.reduce((prev, cur) => prev + cur, 0); if (result < tolerance) { break; } } device.submitPass(computePass); const out = await readback.readBuffer(rBuffer, 0, new Float32Array(r)); invalidation.then(() => { storeProgram.destroy(); storePipeline.destroy(); matmulProgram.destroy(); matmulPipeline.destroy(); rankDiffProgram.destroy(); rankDiffPipeline.destroy(); device.destroy(); // For debug. device.checkForLeaks(); }); return out; })()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:[] 2f:{"id":"da8420861f120e83","slug":"page-rank","title":"Page Rank","thumbnail":"3076384a331419772166c1d7a33ecda4964552d633870eef318301f100ea4e2f","access_level":"public","create_time":"$D2023-12-22T10:17:48.453Z","update_time":"$D2023-12-23T11:55:33.220Z","view_time":null,"first_public_time":"$D2023-12-22T10:17:48.453Z","fork":{"id":"2bd1c0668c8b80c8","slug":"compute-boids","title":"Compute Boids","owner":{"id":"96a4168c70a60e69","login":"antv","name":"AntV"}},"stars":2,"starred":false,"owner":{"type":"team","id":"96a4168c70a60e69","login":"antv","name":"AntV","avatar_url":"https://avatars-next.observableusercontent.com/9cfb1b3742d5fee0ae1105c158818aad1b91fd02751860be879c813cbaac3abf","tier":"starter_2024"},"version":235,"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":2,"collection":null,"type":"notebook","body":{"title":"Page Rank","cells":[{"id":0,"value":"# Page Rank\n\nMigrate the CUDA algorithm to WebGPU.\nhttps://github.com/princeofpython/PageRank-with-CUDA/blob/main/parallel.cu\n\nThe algorithm tries to find the most important node in the graph (all nodes in order of importance).","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\n\nhttps://github.com/princeofpython/PageRank-with-CUDA\n\nWe use 3 kernels here:\n\n```cuda\nwhile(max_iter--){\n\n store_rank<<>>(gpu_r, gpu_r_last, n);\n cudaDeviceSynchronize();\n\n matmul<<>>(gpu_graph, gpu_r, gpu_r_last, n);\n cudaDeviceSynchronize();\n \n rank_diff<<>>(gpu_r, gpu_r_last, n);\n cudaDeviceSynchronize();\n\n cudaMemcpy(r_last, gpu_r_last, n* sizeof(float), cudaMemcpyDeviceToHost);\n float result = thrust::reduce( r_last, r_last + n);\n\n if(result < eps){\n cudaMemcpy(r, gpu_r, n* sizeof(float), cudaMemcpyDeviceToHost);\n return;\n }\n}\ncudaMemcpy(r, gpu_r, n* sizeof(float), cudaMemcpyDeviceToHost);\n```","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":214,"value":"## Kernel Code\n\nWe use 3 storage buffers:\n* r\n* r_last\n* graph\n\nThe store kernel just copy the data from `r` to `r_last`:\n\n```wgsl\nstruct Buffer {\n data: array,\n};\n\n@group(2) @binding(0) var r : Buffer;\n@group(2) @binding(1) var r_last : Buffer;\n\n@compute @workgroup_size(\\${BLOCKS}, \\${BLOCK_SIZE})\nfn main(\n @builtin(global_invocation_id) global_id : vec3\n) {\n var index = global_id.x;\n if (index < \\${V.length}u) {\n r_last.data[index] = r.data[index];\n }\n}\n```","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":232,"value":"The matmul kernel:","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":222,"value":"```wgsl\n@group(0) @binding(0) var graph : Buffer;\n@group(0) @binding(1) var r : Buffer;\n@group(0) @binding(2) var r_last : Buffer;\n\n@compute @workgroup_size(\\${BLOCKS}, \\${BLOCK_SIZE})\nfn main(\n @builtin(global_invocation_id) global_id : vec3\n) {\n var index = global_id.x;\n if (index < \\${V.length}u) {\n var sum = 0.0;\n for (var i = 0u; i < \\${V.length}u; i = i + 1u) {\n sum = sum + r_last.data[i] * graph.data[index * \\${V.length}u + i];\n }\n r.data[index] = sum;\n }\n}\n```","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":234,"value":"The rankDiff kernel:","pinned":false,"hidden":false,"mode":"md","author":"$undefined","summary":"$undefined"},{"id":229,"value":"```wgsl\n@group(0) @binding(0) var r : Buffer;\n@group(0) @binding(1) var r_last : Buffer;\n\n@compute @workgroup_size(\\${BLOCKS}, \\${BLOCK_SIZE})\nfn main(\n @builtin(global_invocation_id) global_id : vec3\n) {\n var index = global_id.x;\n if (index < \\${V.length}u) {\n r_last.data[index] = abs(r_last.data[index] - r.data[index]);\n }\n}\n```","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":224,"value":"1k nodes 50k edges, CPU 13641.50 ms\tGPU 130.20 ms\tspeed up ~100x","pinned":false,"hidden":false,"mode":"md","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"}} 30:"$2f:body"