31:T611,previews([ { path: "@antv/simple-storage", thumbnail: "3307b08de65e4886246b9fa7c57281d2b52dd0c2c37d34fed34575e5a70e6360", title: "Simple Storage" }, { path: "@antv/simplex-noise", thumbnail: "108b2bb8ba34425b5e523417358772739fb5543dcb9b15edf56615ca393a303a", title: "Simplex Noise" }, { path: "@antv/sdf-2d", thumbnail: "a66e7ee0ffafc20868be61a2bfff1180d48ba9cfb25acab6e8394881c72a304a", title: "SDF 2D" }, { path: "@antv/3d-atomics-rasterizer", thumbnail: "cc569cd1fcd44e62ddd6eac8c9f6a34ff9784c09e7ecaea22b93e99197fd3e8a", title: "3D Atomics Rasterizer" }, { path: "@antv/water-caustics", thumbnail: "a78c977dac2343fa054e7d7b593bb013f998be91b11036397009623898fc84e1", title: "Water Caustics" }, { path: "@antv/compute-boids", thumbnail: "a203b1a8a0213096474105c9f8d67ae80faf21898bccf1872d17f9f39532176b", title: "Compute Boids" }, { path: "@antv/blur-image", thumbnail: "70d4488094e287cedc3059b203f4a05f730f92ba00759b09f0a20e611c3858e9", title: "Blur Image" }, { path: "@antv/origami", thumbnail: "7c2ce019170774d081c7f1b658fd81f2b5650b6f7ec07e3c9da1b8b09fda683d", title: "Origami" }, { path: "@antv/stardust", thumbnail: "715df729d1fd43ebe937bc31933c52751021debf3c786f15c806e0ae8926a5fc", title: "Stardust" }, { path: "@antv/page-rank", title: "Page Rank" }, { path: "@antv/sssp", title: "SSSP(Single Source Shortest Path)" } ])32:T591,### Groups and Bindings In our internal implementation, `@group(x)` in WGSL should obey the following order: - `group(0)` Uniform - eg. `var time : Time;` - `group(1)` Texture & Sampler pair - eg. `var pass_in: texture_2d_array; var bilinear: sampler;` - `group(2)` StorageBuffer - eg. `var atomic_storage : array>;` - `group(3)` StorageTexture - eg. `var screen : texture_storage_2d;` And in each group, binding number starts from 0. For example, if we want to use a texture and sampler pair like this: ```wgsl @group(1) @binding(0) var myTexture : texture_2d; @group(1) @binding(1) var mySampler : sampler; ``` When create bindings, default binding number will be assigned. ```ts device.createBindings({ pipeline: renderPipeline, samplerBindings: [ // @group(1) { texture: myTexture, // binding 0 sampler: mySampler, // binding 1 // textureBinding: 10, // Or assign binding number manually // samplerBinding: 11, // samplerBinding: -1, // -1 means we don't need default sampler for this texture }, ], }); ``` In this way, we basically write shader in raw WGSL syntax: * https://observablehq.com/@antv/compute-boids * https://observablehq.com/@antv/blur-image But sometimes we want to use more advanced features that WGSL won't support such as shader chunks, `#define` directive etc.33:T74c,## Shader Chunks (Optional) It's worth mentioning that I try to enhance [WGSL](https://www.w3.org/TR/WGSL/) with shader chunks. Shader chunks can be combined at compiling time or runtime. glslify and its rollup plugin use the former way. But most web rendering engines choose the latter one such as: Three.js, Babylon.js and Orillusion. Bevy use [naga_oil](https://github.com/bevyengine/naga_oil). I try to compile it to WASM with wasm-pack so that features like [preprocessing](https://github.com/bevyengine/naga_oil#preprocessing) and [imports](https://github.com/bevyengine/naga_oil#imports) will come in handy. You will see the following syntax in the example: ```wgsl #import prelude::{screen, time} // import screen texture @compute @workgroup_size(16, 16) fn main_image(@builtin(global_invocation_id) id: uint3) { let screen_size = uint2(textureDimensions(screen)); // use screen texture let t = time.elapsed; // use time uniform } ``` ### Register Shader Module First, register a shader module like this(from https://observablehq.com/@antv/3d-atomics-rasterizer): ```wgsl #define_import_path custom // module name // export a struct struct Custom { Radius: f32, } // export a var @group(0) @binding(1) var custom: Custom; // export a function fn test() -> uint4 {} ``` Then calling our wrapped method called [registerShaderModule](https://observablehq.com/@antv/compute-toys#registerShaderModule): ```ts registerShaderModule(custom); ``` ### Use Shader Module Now you can import from this shader module like this: ```wgsl #import custom::{custom, test}; custom.Radius; // f32 test(); ``` Don't forget to use our wrapped method called [createProgram](https://observablehq.com/@antv/compute-toys#createProgram) instead of the one from device: ```ts // before device.createProgram({}); // after createProgram(device, {}); ```34:T10f7,## Diving Deep into Compute Toys https://compute.toys/ makes developing Compute Shaders easy. They let developers focus on the Shader itself instead of the API. * Prelude includes alias, variables and helper functions. [Source Code](https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/lib.rs#L367) * Directives like `#define`, `#include`, `#dispatch_count`. [Source Code]() * Input and output. An internal blit renderpass can output computed result to screen. [Source Code](https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/blit.wgsl) How about we mimic the implementation based on what we already know? These examples are written in this way and migrated directly from https://compute.toys/: * https://observablehq.com/@antv/simple-storage * https://observablehq.com/@antv/simplex-noise * https://observablehq.com/@antv/3d-atomics-rasterizer * https://observablehq.com/@antv/water-caustics ### Blit RenderPass Fullscreen shader is really common in postprocessing. Compute Toys uses a [fullscreen triangle](https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/blit.wgsl). [Bevy](https://bevyengine.org/) gives a very detailed explanation: ``` // This vertex shader produces the following, when drawn using indices 0..3: // // 1 | 0-----x.....2 // 0 | | s | . ´ // -1 | x_____x´ // -2 | : .´ // -3 | 1´ // +--------------- // -1 0 1 2 3 // // The axes are clip-space x and y. The region marked s is the visible region. // The digits in the corners of the right-angled triangle are the vertex // indices. // // The top-left has UV 0,0, the bottom-left has 0,2, and the top-right has 2,0. // This means that the UV gets interpolated to 1,1 at the bottom-right corner // of the clip-space rectangle that is at 1,-1 in clip space. ``` We use the same fullscreen triangle in [createBlitPipelineAndBindings](https://observablehq.com/@antv/compute-toys#createBlitPipelineAndBindings), which uses less vertices count compared with a quad. ### Pass in and out The shader can write to pass_out, which will be copied into pass_in after the current entrypoint has returned. https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/lib.rs#L270 In my implementation, when compute pass submitted we can do the same thing with `copySubTexture2D`: ```ts // https://observablehq.com/@antv/water-caustics device.submitPass(computePass); device.copySubTexture2D(pass_in, 0, 0, pass_out, 0, 0, 4); ``` BTW, compute toys create two 128 MB storages. [Source code](https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/bind.rs#L356) ### Alias ```wgsl alias int = i32; alias uint = u32; alias float = f32; alias int2 = vec2; // ... ``` ### Uniform Compute toys provide uniform GUI on the left panel. We create uniform buffer and update on each frame. With the Input on observable, these params can be changed. ```ts const customUniformBuffer = device.createBuffer({ viewOrSize: 11 * Float32Array.BYTES_PER_ELEMENT, usage: BufferUsage.UNIFORM, }); customUniformBuffer.setSubData( 0, new Uint8Array( new Float32Array([ 0.551, 0.313, 0.632, 0.489, 0.018, 0.697, 0.621, 0, 1, 0.962, 0.224, ]).buffer, ), ); ``` ### Prelude With capabilities of [shader chunks](https://observablehq.com/@antv/compute-toys#cell-712), we register the prelude module in every example like this: ```ts registerShaderModule(prelude); ``` The shader block looks the same as the one used in compute toys. The group indices are changed due to our internal rules. ```wgsl #define_import_path prelude struct Time { frame: u32, elapsed : f32 } @group(0) @binding(0) var time : Time; @group(3) @binding(0) var screen : texture_storage_2d; @group(1) @binding(0) var pass_in: texture_2d_array; @group(1) @binding(1) var bilinear: sampler; @group(3) @binding(1) var pass_out: texture_storage_2d_array; fn passStore(pass_index: int, coord: int2, value: float4) { textureStore(pass_out, coord, pass_index, value); } fn passLoad(pass_index: int, coord: int2, lod: int) -> float4 { return textureLoad(pass_in, coord, pass_index, lod); } fn passSampleLevelBilinearRepeat(pass_index: int, uv: float2, lod: float) -> float4 { return textureSampleLevel(pass_in, bilinear, fract(uv), pass_index, lod); } ```35:Tcb5,createBlitPipelineAndBindings = (device, screen) => { const sampler = device.createSampler({ addressModeU: deviceAPI.AddressMode.CLAMP_TO_EDGE, addressModeV: deviceAPI.AddressMode.CLAMP_TO_EDGE, minFilter: deviceAPI.FilterMode.BILINEAR, magFilter: deviceAPI.FilterMode.BILINEAR, mipmapFilter: deviceAPI.MipmapFilterMode.LINEAR }); const renderProgram = device.createProgram({ vertex: { entryPoint: "fullscreen_vertex_shader", wgsl: ` struct FullscreenVertexOutput { @builtin(position) position: vec4, @location(0) uv: vec2, }; // This vertex shader produces the following, when drawn using indices 0..3: // // 1 | 0-----x.....2 // 0 | | s | . ´ // -1 | x_____x´ // -2 | : .´ // -3 | 1´ // +--------------- // -1 0 1 2 3 // // The axes are clip-space x and y. The region marked s is the visible region. // The digits in the corners of the right-angled triangle are the vertex // indices. // // The top-left has UV 0,0, the bottom-left has 0,2, and the top-right has 2,0. // This means that the UV gets interpolated to 1,1 at the bottom-right corner // of the clip-space rectangle that is at 1,-1 in clip space. @vertex fn fullscreen_vertex_shader(@builtin(vertex_index) vertex_index: u32) -> FullscreenVertexOutput { // See the explanation above for how this works let uv = vec2(f32(vertex_index >> 1u), f32(vertex_index & 1u)) * 2.0; let clip_position = vec4(uv * vec2(2.0, -2.0) + vec2(-1.0, 1.0), 0.0, 1.0); return FullscreenVertexOutput(clip_position, uv); } ` }, fragment: { entryPoint: "fs_main_linear_to_srgb", wgsl: ` struct FullscreenVertexOutput { @builtin(position) position: vec4, @location(0) uv: vec2, }; @group(1) @binding(0) var in_texture : texture_2d; @group(1) @binding(1) var in_sampler : sampler; fn srgb_to_linear(rgb: vec3) -> vec3 { return select( pow((rgb + 0.055) * (1.0 / 1.055), vec3(2.4)), rgb * (1.0/12.92), rgb <= vec3(0.04045)); } fn linear_to_srgb(rgb: vec3) -> vec3 { return select( 1.055 * pow(rgb, vec3(1.0 / 2.4)) - 0.055, rgb * 12.92, rgb <= vec3(0.0031308)); } @fragment fn fs_main(in: FullscreenVertexOutput) -> @location(0) vec4 { return textureSample(in_texture, in_sampler, in.uv); } @fragment fn fs_main_linear_to_srgb(in: FullscreenVertexOutput) -> @location(0) vec4 { let rgba = textureSample(in_texture, in_sampler, in.uv); return vec4(linear_to_srgb(rgba.rgb), rgba.a); } @fragment fn fs_main_rgbe_to_linear(in: FullscreenVertexOutput) -> @location(0) vec4 { let rgbe = textureSample(in_texture, in_sampler, in.uv); return vec4(rgbe.rgb * exp2(rgbe.a * 255. - 128.), 1.); } ` } }); const renderPipeline = device.createRenderPipeline({ inputLayout: null, program: renderProgram, colorAttachmentFormats: [deviceAPI.Format.U8_RGBA_RT] }); const showResultBindings = device.createBindings({ pipeline: renderPipeline, samplerBindings: [ { texture: screen, sampler } ] }); return { pipeline: renderPipeline, bindings: showResultBindings }; }36:[] 2f:{"title":"Compute Toys (with @antv/g-device-api@1.6.4)","cells":[{"id":0,"value":"# Compute Toys (with @antv/g-device-api@1.6.4)","pinned":false,"hidden":false,"mode":"md"},{"id":67,"value":"A collection of WebGPU Compute Shaders, inspired by https://compute.toys/\n\nI used my HAL called [@antv/g-device-api](https://observablehq.com/@antv/g-device-api) underlying. However, since WebGL does not support Compute Shader at all, you can use native WebGPU API directly.\n\n* [Getting Started](#cell-32)\n* [Shader Chunks](#cell-712)\n* [Diving Deep into Compute Toys](#cell-794)\n\nIn our scenario, we will try to use GPGPU in graph algorithms and music visualization.\n\nHere are some useful references of Compute Shaders:\n* https://compute.toys/\n* https://webgpufundamentals.org/webgpu/lessons/webgpu-compute-shaders.html\n* https://webgpufundamentals.org/webgpu/lessons/webgpu-compute-shaders-histogram.html","pinned":false,"hidden":false,"mode":"md"},{"id":675,"value":"$31","pinned":false,"hidden":false,"mode":"ojs"},{"id":32,"value":"## Getting Started\n\n### Create WebGPU Device\n\nCreate a WebGPU Device contribution.\n```ts\nimport { WebGPUDeviceContribution } from '@antv/g-device-api';\nconst deviceContribution = new WebGPUDeviceContribution({\n shaderCompilerPath: '/glsl_wgsl_compiler_bg.wasm',\n // From CDN\n // shaderCompilerPath: 'https://unpkg.com/@antv/g-device-api@1.4.9/rust/pkg/glsl_wgsl_compiler_bg.wasm',\n});\n```","pinned":false,"hidden":false,"mode":"md"},{"id":179,"value":"Then create swapChain with contribution and get device from it:\n```ts\nconst swapChain = await deviceContribution.createSwapChain($canvas);\nswapChain.configureSwapChain($canvas.width, $canvas.height);\nconst device = swapChain.getDevice();\n```","pinned":false,"hidden":false,"mode":"md"},{"id":183,"value":"Now we can use [Device API](https://www.npmjs.com/package/@strawberry-vis/g-device-api?activeTab=readme#api) with the created device.\n\n### Create Compute Pipeline\n\n```ts\nconst computeProgram = device.createProgram({\n compute: {\n entryPoint: 'main',\n wgsl: `` // Raw WGSL\n }\n});\n```\n\nThen we need to create a ComputePipeline and link to compute program:\n\n```ts\nconst computePipeline = device.createComputePipeline({\n inputLayout: null,\n program: computeProgram,\n});\n```\n\nOn each frame, a new ComputePass should be created and submitted at the end:\n\n```ts\nconst computePass = device.createComputePass();\ncomputePass.setPipeline(computePipeline);\ncomputePass.setBindings(bindings);\ncomputePass.dispatchWorkgroups(1);\ndevice.submitPass(computePass);\n```\n\nBTW, the groups and bindings may be a little bit messy now.","pinned":false,"hidden":false,"mode":"md"},{"id":828,"value":"$32","pinned":false,"hidden":false,"mode":"md"},{"id":712,"value":"$33","pinned":false,"hidden":false,"mode":"md"},{"id":774,"value":"### Preprocessor Directives","pinned":false,"hidden":false,"mode":"md"},{"id":792,"value":"Conditional compilation is performed with `#if` / `#ifdef` / `#ifndef`, `#else` and `#endif` preprocessor directives:\n\n```wgsl\nfn get_number() -> f32 {\n #ifdef BIG_NUMBER\n return 999.0;\n #else\n return 0.999;\n #endif\n}\n```","pinned":false,"hidden":false,"mode":"md"},{"id":794,"value":"$34","pinned":false,"hidden":false,"mode":"md"},{"id":832,"value":"### Limitations\n\nIf the property inside struct ends with number, naga-oit will throw an error when compiling.\n\n```wgsl\n// error\nstruct Custom {\n Sine1: f32,\n}\n\n// correct\nstruct Custom {\n SineA: f32,\n}\n```\n\nFor now we create storage texture in `F16_RGBA` format:\n\n```ts\nconst screen = device.createTexture({\n format: Format.F16_RGBA, // Format.F32_RGBA will throw an error\n width: $canvas.width,\n height: $canvas.height,\n dimension: TextureDimension.TEXTURE_2D,\n usage: TextureUsage.STORAGE,\n});\n```\n\n- https://www.w3.org/TR/webgpu/#float32-filterable\n- https://github.com/compute-toys/wgpu-compute-toy/blob/master/src/bind.rs#L433","pinned":false,"hidden":false,"mode":"md"},{"id":40,"value":"## Appendix","pinned":false,"hidden":false,"mode":"md"},{"id":671,"value":"import { previews } from \"@d3/gallery\"","pinned":false,"hidden":false,"mode":"ojs"},{"id":725,"value":"deviceAPI = require(\"@antv/g-device-api@1.6.4\")","pinned":false,"hidden":false,"mode":"ojs"},{"id":731,"value":"createCanvas = () => {\n const $canvas = document.createElement(\"canvas\");\n $canvas.width = 1000;\n $canvas.height = 1000;\n $canvas.style.width = \"500px\";\n $canvas.style.height = \"500px\";\n return $canvas;\n}","pinned":false,"hidden":false,"mode":"ojs"},{"id":728,"value":"createDeviceContributionWebGPU = () =>\n new deviceAPI.WebGPUDeviceContribution({\n shaderCompilerPath:\n \"https://unpkg.com/@antv/g-device-api@1.6.4/rust/pkg/glsl_wgsl_compiler_bg.wasm\"\n })","pinned":true,"hidden":false,"mode":"ojs"},{"id":746,"value":"createProgram = (device, desc, defines) => {\n const compiler = device[\"WGSLComposer\"];\n\n // Prepend defines.\n const prefix =\n Object.keys(defines || {})\n .map((key) => {\n return `#define ${key}`;\n })\n .join(\"\\n\") + \"\\n\";\n\n Object.keys(desc).forEach((key) => {\n desc[key].wgsl = alias + desc[key].wgsl;\n\n if (desc[key].defines) {\n desc[key].wgsl = prefix + desc[key].wgsl;\n }\n // Use naga-oil to combine shaders.\n desc[key].wgsl = compiler.wgsl_compile(desc[key].wgsl);\n });\n\n return device.createProgram(desc);\n}","pinned":false,"hidden":false,"mode":"ojs"},{"id":734,"value":"$35","pinned":false,"hidden":false,"mode":"ojs"},{"id":737,"value":"registerShaderModule = (device, shader) => {\n const compiler = device[\"WGSLComposer\"];\n return compiler.wgsl_compile(alias + shader);\n}","pinned":false,"hidden":false,"mode":"ojs"},{"id":739,"value":"alias = `\nalias int = i32;\nalias uint = u32;\nalias float = f32;\nalias int2 = vec2;\nalias int3 = vec3;\nalias int4 = vec4;\nalias uint2 = vec2;\nalias uint3 = vec3;\nalias uint4 = vec4;\nalias float2 = vec2;\nalias float3 = vec3;\nalias float4 = vec4;\nalias bool2 = vec2;\nalias bool3 = vec3;\nalias bool4 = vec4;\nalias float2x2 = mat2x2;\nalias float2x3 = mat2x3;\nalias float2x4 = mat2x4;\nalias float3x2 = mat3x2;\nalias float3x3 = mat3x3;\nalias float3x4 = mat3x4;\nalias float4x2 = mat4x2;\nalias float4x3 = mat4x3;\nalias float4x4 = mat4x4;\n`","pinned":false,"hidden":false,"mode":"ojs"},{"id":742,"value":"prelude = `\n#define_import_path prelude\n\nstruct Time {\n frame: u32,\n elapsed : f32\n}\n\n@group(0) @binding(0) var time : Time;\n@group(3) @binding(0) var screen : texture_storage_2d;\n\n@group(1) @binding(0) var pass_in: texture_2d_array;\n@group(1) @binding(1) var bilinear: sampler;\n@group(3) @binding(1) var pass_out: texture_storage_2d_array;\n\nfn passStore(pass_index: int, coord: int2, value: float4) {\n textureStore(pass_out, coord, pass_index, value);\n}\n\nfn passLoad(pass_index: int, coord: int2, lod: int) -> float4 {\n return textureLoad(pass_in, coord, pass_index, lod);\n}\n\nfn passSampleLevelBilinearRepeat(pass_index: int, uv: float2, lod: float) -> float4 {\n return textureSampleLevel(pass_in, bilinear, fract(uv), pass_index, lod);\n}\n`","pinned":false,"hidden":false,"mode":"ojs"},{"id":518,"value":"glMatrix = require(\"gl-matrix@3.4.2/dist/gl-matrix.js\")","pinned":false,"hidden":false,"mode":"ojs"},{"id":554,"value":"loadImage = async (url) => {\n if (!!window.createImageBitmap) {\n const response = await fetch(url);\n const imageBitmap = await createImageBitmap(await response.blob());\n return imageBitmap;\n } else {\n const image = new window.Image();\n return new Promise((res) => {\n image.onload = () => res(image);\n image.src = url;\n image.crossOrigin = \"Anonymous\";\n });\n }\n}","pinned":false,"hidden":false,"mode":"ojs"}],"files":"$W36","stdlib":"1","titling":"auto"}