diff --git a/.github/workflows/static.yml b/.github/workflows/static.yml index 9432148..46021d2 100644 --- a/.github/workflows/static.yml +++ b/.github/workflows/static.yml @@ -33,7 +33,7 @@ jobs: - name: Set up Node uses: actions/setup-node@v4 with: - node-version: 20 + node-version: 22 cache: 'npm' - name: Install dependencies run: npm ci diff --git a/.gitignore b/.gitignore index 054e565..af94af9 100644 --- a/.gitignore +++ b/.gitignore @@ -24,4 +24,4 @@ dist-ssr *.sw? /.vite -*/scenes \ No newline at end of file +scenes \ No newline at end of file diff --git a/README.md b/README.md index edffdaf..3914726 100644 --- a/README.md +++ b/README.md @@ -1,26 +1,155 @@ -# Project5-WebGPU-Gaussian-Splat-Viewer +# WebGPU Gaussian Splatting +3D Gaussian splatting renderer and viewer in WebGPU. -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 5** +**Author:** Lu M. +- [LinkedIn](https://www.linkedin.com/in/lu-m-673425323/) +- [Personal Site](lu-m-dev.github.io) -* (TODO) YOUR NAME HERE -* Tested on: (TODO) **Google Chrome 222.2** on - Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +**Tested System:** + - Windows 11 Home + - AMD Ryzen 7 5800HS @ 3.20GHz, 16GB RAM + - NVIDIA GeForce RTX 3060 Laptop GPU 6GB (Compute Capability 8.6) + - Brave 1.83.118 (Official Build) (64-bit), Chromium: 141.0.7390.108 -### Live Demo +## Abstract -[![](img/thumb.png)](http://TODO.github.io/Project4-WebGPU-Forward-Plus-and-Clustered-Deferred) +Gaussian splatting is a point-based rendering technique where each 3D point is rasterized as a screen-space Gaussian instead of a single pixel. Each splat carries a position, covariance, and color. When splats overlap, their contributions are composited to reconstruct a smooth, continuous appearance. -### Demo Video/GIF +## Methods -[![](img/video.mp4)](TODO) +This implementation follows three main stages: preprocess, sort, and render. Each stage is WebGPU-parallelized. -### (TODO: Your README) +### Preprocess -*DO NOT* leave the README to the last minute! It is a crucial part of the -project, and we will not be able to grade you without a good README. +- Input a 3D point-cloud loaded from a PLY. +- Transform points into view, clip and NDC coordinates. +- Compute per-point splat parameters: positions, 3D covariance, 2D covariance. -This assignment has a considerable amount of performance analysis compared -to implementation work. Complete the implementation early to leave time! +### Sort + +- Purpose: sorting points back-to-front is required for correct alpha blending and improved memory coherence. +- Implementation: a GPU radix sort. + +### Render + +Two renderers are provided: +- Point-cloud renderer: renders raw points with a simple size and per-point color. This uses a minimal vertex/fragment pipeline. +- Gaussian splat renderer: for each splat, a screen-space quad is rasterized and shaded with a precomputed Gaussian weight using the splat covariance and color. The fragment shader evaluates the color and opacity. + +### Tile-Based Depth Sorting +Improve rendering performance on large point clouds by organizing Gaussian splats into screen-space tiles and sorting per-tile, reducing memory bandwidth and improving cache locality. + +Implementation: a compute-based pipeline that: +1. For each Gaussian, determines which tiles it overlaps using its 2D projection and splat radius. +2. Generates key-value pairs `(tile_id | depth, gaussian_id)` for each Gaussian-tile intersection using atomic operations to avoid write conflicts. +3. Sorts these pairs globally using radix sort. +4. Identifies start/end ranges for each tile via a second compute pass. +5. During rendering, tiles can be rasterized in any order with per-tile depth ordering. + +Trade-offs: This approach reduces per-frame fragment shader overdraw in theory but adds preprocessing overhead. For scenes with significant depth complexity and overlapping splats per tile, this can improve performance. For scenes with sparse or well-distributed splats, the overhead may outweigh the benefits. + +## Visual comparison + +### Bonsai (272,965 points) + + + + + + + + + + + + + + +
Point Cloud RendererGaussian Splatting
bonsai pointcloudbonsai gaussian
272,965 points — point renderer272,965 points — gaussian splatting
+ +Observation: for the bonsai dataset (272,965 points) both renderers saturate the display refresh on the test laptop and hit the monitor's 144 Hz cap in the author's measurements. + +### Bicycle (1,063,091 points) + + + + + + + + + + + + + + +
Point Cloud RendererGaussian Splatting
bicycle pointcloudbicycle gaussian
1,063,091 points — point renderer1,063,091 points — gaussian splatting
+ +Measured performance (personal laptop): + +- Bonsai (272,965 points): both renderers hit the display cap at 144 Hz. +- Bicycle (1,063,091 points): point-cloud renderer ≈ 100 Hz; Gaussian splatting renderer ≈ 50 Hz; Gaussian splatting with tile-based depth sorting ≈ 40 Hz. + +Discussion: the Gaussian splatting renderer rasterizes a screen-space quad per point and evaluates a Gaussian per-fragment. For dense clouds (the bicycle set) this produces significantly more fragment shader work and overdraw compared to rendering simple points. For the bonsai dataset the total fragment workload is low enough that both appear similarly fast. + +The tile-based depth sorting variant demonstrates no visual difference in rendered quality compared to the base Gaussian splatting renderer. However, it runs approximately 20% slower (40 FPS vs 50 FPS on bicycle) despite the theoretical benefits of improved memory coherence and reduced overdraw. This performance regression occurs because: +1. **Preprocessing overhead dominates**: the tile-based approach adds significant compute cost in key generation, atomic allocations, sorting the larger tile-pair dataset, and range identification. These stages are necessary every frame and account for substantial GPU time. +2. **Tile-sorting benefits don't materialize at this scale**: while tile-based sorting theoretically improves cache locality, the actual savings in fragment shader evaluation are minimal. The 1M point cloud is already sparse enough per tile that depth coherence doesn't yield meaningful performance gains on modern GPUs with efficient caching. +3. **Atomic operation contention**: multiple Gaussians writing to the same tile allocator creates atomic contention, serializing work that could otherwise be parallelized. +4. **Small tile sizes**: to fit within WebGPU buffer limits (max_tile_pairs ≈ 16M), tile sizes must be modest, spreading splats across many tiles and reducing the per-tile depth-sorting benefit. + +Note: these are preliminary numbers from a single machine. More comprehensive benchmarking is required (different GPUs, tile-based profiling, memory bandwidth counters, and varying splat sizes) to draw robust conclusions. + +## Bloopers +### Incorrect depth and sort buffer +
+bonsai blooper +
+I implemented incorrect depth calculation and improper sort buffer usage. This resulted in the output displaying Gaussian quads refreshing with black square artifacts that obscured the view. + +### Mistaken opacity + + + + + + + + + +
BonsaiBicycle
bonsai blooperbicycle blooper
+I implemented wrong alpha calculation in the fragment shader, mistakenly assigning an arbitrary value of `1.0f` to all fragments. This gave the air around the object an opacity coherent with the object's color. + +### Tile-based sorting buffer overflow +
+tile sorting blooper +
+During tile-based depth sorting, I failed to properly allocate space for tile-pair entries using atomic operations, causing buffer overwrites. This resulted in splats rendering as solid colored discs rather than smoothly blended Gaussians, due to corrupted splat data and invalid depth values in the tile-sorted output. + +## Build instructions + +1) Download [Node.js](https://nodejs.org/en/download) + +2) Clone repository + ```cmd + git clone https://github.com/lu-m-dev/WebGPU-gaussian-splatting.git + ``` + +3) Install dependencies + ```cmd + cd WebGPU-gaussian-splatting + npm install + ``` + +3) Launch dev server + ```cmd + npm run dev + ``` + Optional build to static `dist/` + ```cmd + npm run build + ``` ### Credits diff --git a/images/bicycle_gaussian.png b/images/bicycle_gaussian.png new file mode 100644 index 0000000..7d2a7f6 Binary files /dev/null and b/images/bicycle_gaussian.png differ diff --git a/images/bicycle_pointcloud.png b/images/bicycle_pointcloud.png new file mode 100644 index 0000000..9d56a5f Binary files /dev/null and b/images/bicycle_pointcloud.png differ diff --git a/images/blooper.gif b/images/blooper.gif new file mode 100644 index 0000000..ae0bcea Binary files /dev/null and b/images/blooper.gif differ diff --git a/images/blooper_bicycle.png b/images/blooper_bicycle.png new file mode 100644 index 0000000..a112558 Binary files /dev/null and b/images/blooper_bicycle.png differ diff --git a/images/blooper_bonsai.png b/images/blooper_bonsai.png new file mode 100644 index 0000000..ccf58ac Binary files /dev/null and b/images/blooper_bonsai.png differ diff --git a/images/blooper_tile.gif b/images/blooper_tile.gif new file mode 100644 index 0000000..5c8be9a Binary files /dev/null and b/images/blooper_tile.gif differ diff --git a/images/bonsai_gaussian.png b/images/bonsai_gaussian.png new file mode 100644 index 0000000..550d364 Binary files /dev/null and b/images/bonsai_gaussian.png differ diff --git a/images/bonsai_pointcloud.png b/images/bonsai_pointcloud.png new file mode 100644 index 0000000..cf300d6 Binary files /dev/null and b/images/bonsai_pointcloud.png differ diff --git a/images/cover.gif b/images/cover.gif new file mode 100644 index 0000000..d75d982 Binary files /dev/null and b/images/cover.gif differ diff --git a/images/preview.gif b/images/preview.gif new file mode 100644 index 0000000..fde2c2d Binary files /dev/null and b/images/preview.gif differ diff --git a/index.html b/index.html index 4067a69..674fa65 100644 --- a/index.html +++ b/index.html @@ -3,7 +3,7 @@ - WebGPU 3D Gaussian Splat Viewer + WebGPU-gaussian-splatting
diff --git a/package-lock.json b/package-lock.json index 04843bd..2e2f09d 100644 --- a/package-lock.json +++ b/package-lock.json @@ -12,6 +12,7 @@ "@loaders.gl/ply": "^4.2.2", "@petamoriken/float16": "^3.8.7", "tweakpane": "^3.1.8", + "tweakpane-plugin-file-import": "^0.2.0", "wgpu-matrix": "^3.2.0" }, "devDependencies": { @@ -19,366 +20,457 @@ "@webgpu/types": "^0.1.31", "tweakpane-plugin-file-import": "^0.2.0", "typescript": "^5.0.4", - "vite": "^4.3.1", + "vite": "^7.1.12", "vite-raw-plugin": "^1.0.1" } }, + "node_modules/@esbuild/aix-ppc64": { + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/aix-ppc64/-/aix-ppc64-0.25.11.tgz", + "integrity": "sha512-Xt1dOL13m8u0WE8iplx9Ibbm+hFAO0GsU2P34UNoDGvZYkY8ifSiy6Zuc1lYxfG7svWE2fzqCUmFp5HCn51gJg==", + "cpu": [ + "ppc64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "aix" + ], + "engines": { + "node": ">=18" + } + }, "node_modules/@esbuild/android-arm": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/android-arm/-/android-arm-0.18.20.tgz", - "integrity": "sha512-fyi7TDI/ijKKNZTUJAQqiG5T7YjJXgnzkURqmGj13C6dCqckZBLdl4h7bkhHt/t0WP+zO9/zwroDvANaOqO5Sw==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/android-arm/-/android-arm-0.25.11.tgz", + "integrity": "sha512-uoa7dU+Dt3HYsethkJ1k6Z9YdcHjTrSb5NUy66ZfZaSV8hEYGD5ZHbEMXnqLFlbBflLsl89Zke7CAdDJ4JI+Gg==", "cpu": [ "arm" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "android" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/android-arm64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/android-arm64/-/android-arm64-0.18.20.tgz", - "integrity": "sha512-Nz4rJcchGDtENV0eMKUNa6L12zz2zBDXuhj/Vjh18zGqB44Bi7MBMSXjgunJgjRhCmKOjnPuZp4Mb6OKqtMHLQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/android-arm64/-/android-arm64-0.25.11.tgz", + "integrity": "sha512-9slpyFBc4FPPz48+f6jyiXOx/Y4v34TUeDDXJpZqAWQn/08lKGeD8aDp9TMn9jDz2CiEuHwfhRmGBvpnd/PWIQ==", "cpu": [ "arm64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "android" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/android-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/android-x64/-/android-x64-0.18.20.tgz", - "integrity": "sha512-8GDdlePJA8D6zlZYJV/jnrRAi6rOiNaCC/JclcXpB+KIuvfBN4owLtgzY2bsxnx666XjJx2kDPUmnTtR8qKQUg==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/android-x64/-/android-x64-0.25.11.tgz", + "integrity": "sha512-Sgiab4xBjPU1QoPEIqS3Xx+R2lezu0LKIEcYe6pftr56PqPygbB7+szVnzoShbx64MUupqoE0KyRlN7gezbl8g==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "android" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/darwin-arm64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/darwin-arm64/-/darwin-arm64-0.18.20.tgz", - "integrity": "sha512-bxRHW5kHU38zS2lPTPOyuyTm+S+eobPUnTNkdJEfAddYgEcll4xkT8DB9d2008DtTbl7uJag2HuE5NZAZgnNEA==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/darwin-arm64/-/darwin-arm64-0.25.11.tgz", + "integrity": "sha512-VekY0PBCukppoQrycFxUqkCojnTQhdec0vevUL/EDOCnXd9LKWqD/bHwMPzigIJXPhC59Vd1WFIL57SKs2mg4w==", "cpu": [ "arm64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "darwin" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/darwin-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/darwin-x64/-/darwin-x64-0.18.20.tgz", - "integrity": "sha512-pc5gxlMDxzm513qPGbCbDukOdsGtKhfxD1zJKXjCCcU7ju50O7MeAZ8c4krSJcOIJGFR+qx21yMMVYwiQvyTyQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/darwin-x64/-/darwin-x64-0.25.11.tgz", + "integrity": "sha512-+hfp3yfBalNEpTGp9loYgbknjR695HkqtY3d3/JjSRUyPg/xd6q+mQqIb5qdywnDxRZykIHs3axEqU6l1+oWEQ==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "darwin" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/freebsd-arm64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/freebsd-arm64/-/freebsd-arm64-0.18.20.tgz", - "integrity": "sha512-yqDQHy4QHevpMAaxhhIwYPMv1NECwOvIpGCZkECn8w2WFHXjEwrBn3CeNIYsibZ/iZEUemj++M26W3cNR5h+Tw==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/freebsd-arm64/-/freebsd-arm64-0.25.11.tgz", + "integrity": "sha512-CmKjrnayyTJF2eVuO//uSjl/K3KsMIeYeyN7FyDBjsR3lnSJHaXlVoAK8DZa7lXWChbuOk7NjAc7ygAwrnPBhA==", "cpu": [ "arm64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "freebsd" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/freebsd-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/freebsd-x64/-/freebsd-x64-0.18.20.tgz", - "integrity": "sha512-tgWRPPuQsd3RmBZwarGVHZQvtzfEBOreNuxEMKFcd5DaDn2PbBxfwLcj4+aenoh7ctXcbXmOQIn8HI6mCSw5MQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/freebsd-x64/-/freebsd-x64-0.25.11.tgz", + "integrity": "sha512-Dyq+5oscTJvMaYPvW3x3FLpi2+gSZTCE/1ffdwuM6G1ARang/mb3jvjxs0mw6n3Lsw84ocfo9CrNMqc5lTfGOw==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "freebsd" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-arm": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-arm/-/linux-arm-0.18.20.tgz", - "integrity": "sha512-/5bHkMWnq1EgKr1V+Ybz3s1hWXok7mDFUMQ4cG10AfW3wL02PSZi5kFpYKrptDsgb2WAJIvRcDm+qIvXf/apvg==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-arm/-/linux-arm-0.25.11.tgz", + "integrity": "sha512-TBMv6B4kCfrGJ8cUPo7vd6NECZH/8hPpBHHlYI3qzoYFvWu2AdTvZNuU/7hsbKWqu/COU7NIK12dHAAqBLLXgw==", "cpu": [ "arm" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-arm64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-arm64/-/linux-arm64-0.18.20.tgz", - "integrity": "sha512-2YbscF+UL7SQAVIpnWvYwM+3LskyDmPhe31pE7/aoTMFKKzIc9lLbyGUpmmb8a8AixOL61sQ/mFh3jEjHYFvdA==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-arm64/-/linux-arm64-0.25.11.tgz", + "integrity": "sha512-Qr8AzcplUhGvdyUF08A1kHU3Vr2O88xxP0Tm8GcdVOUm25XYcMPp2YqSVHbLuXzYQMf9Bh/iKx7YPqECs6ffLA==", "cpu": [ "arm64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-ia32": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-ia32/-/linux-ia32-0.18.20.tgz", - "integrity": "sha512-P4etWwq6IsReT0E1KHU40bOnzMHoH73aXp96Fs8TIT6z9Hu8G6+0SHSw9i2isWrD2nbx2qo5yUqACgdfVGx7TA==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-ia32/-/linux-ia32-0.25.11.tgz", + "integrity": "sha512-TmnJg8BMGPehs5JKrCLqyWTVAvielc615jbkOirATQvWWB1NMXY77oLMzsUjRLa0+ngecEmDGqt5jiDC6bfvOw==", "cpu": [ "ia32" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-loong64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-loong64/-/linux-loong64-0.18.20.tgz", - "integrity": "sha512-nXW8nqBTrOpDLPgPY9uV+/1DjxoQ7DoB2N8eocyq8I9XuqJ7BiAMDMf9n1xZM9TgW0J8zrquIb/A7s3BJv7rjg==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-loong64/-/linux-loong64-0.25.11.tgz", + "integrity": "sha512-DIGXL2+gvDaXlaq8xruNXUJdT5tF+SBbJQKbWy/0J7OhU8gOHOzKmGIlfTTl6nHaCOoipxQbuJi7O++ldrxgMw==", "cpu": [ "loong64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-mips64el": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-mips64el/-/linux-mips64el-0.18.20.tgz", - "integrity": "sha512-d5NeaXZcHp8PzYy5VnXV3VSd2D328Zb+9dEq5HE6bw6+N86JVPExrA6O68OPwobntbNJ0pzCpUFZTo3w0GyetQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-mips64el/-/linux-mips64el-0.25.11.tgz", + "integrity": "sha512-Osx1nALUJu4pU43o9OyjSCXokFkFbyzjXb6VhGIJZQ5JZi8ylCQ9/LFagolPsHtgw6himDSyb5ETSfmp4rpiKQ==", "cpu": [ "mips64el" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-ppc64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-ppc64/-/linux-ppc64-0.18.20.tgz", - "integrity": "sha512-WHPyeScRNcmANnLQkq6AfyXRFr5D6N2sKgkFo2FqguP44Nw2eyDlbTdZwd9GYk98DZG9QItIiTlFLHJHjxP3FA==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-ppc64/-/linux-ppc64-0.25.11.tgz", + "integrity": "sha512-nbLFgsQQEsBa8XSgSTSlrnBSrpoWh7ioFDUmwo158gIm5NNP+17IYmNWzaIzWmgCxq56vfr34xGkOcZ7jX6CPw==", "cpu": [ "ppc64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-riscv64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-riscv64/-/linux-riscv64-0.18.20.tgz", - "integrity": "sha512-WSxo6h5ecI5XH34KC7w5veNnKkju3zBRLEQNY7mv5mtBmrP/MjNBCAlsM2u5hDBlS3NGcTQpoBvRzqBcRtpq1A==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-riscv64/-/linux-riscv64-0.25.11.tgz", + "integrity": "sha512-HfyAmqZi9uBAbgKYP1yGuI7tSREXwIb438q0nqvlpxAOs3XnZ8RsisRfmVsgV486NdjD7Mw2UrFSw51lzUk1ww==", "cpu": [ "riscv64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-s390x": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-s390x/-/linux-s390x-0.18.20.tgz", - "integrity": "sha512-+8231GMs3mAEth6Ja1iK0a1sQ3ohfcpzpRLH8uuc5/KVDFneH6jtAJLFGafpzpMRO6DzJ6AvXKze9LfFMrIHVQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-s390x/-/linux-s390x-0.25.11.tgz", + "integrity": "sha512-HjLqVgSSYnVXRisyfmzsH6mXqyvj0SA7pG5g+9W7ESgwA70AXYNpfKBqh1KbTxmQVaYxpzA/SvlB9oclGPbApw==", "cpu": [ "s390x" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/linux-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/linux-x64/-/linux-x64-0.18.20.tgz", - "integrity": "sha512-UYqiqemphJcNsFEskc73jQ7B9jgwjWrSayxawS6UVFZGWrAAtkzjxSqnoclCXxWtfwLdzU+vTpcNYhpn43uP1w==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/linux-x64/-/linux-x64-0.25.11.tgz", + "integrity": "sha512-HSFAT4+WYjIhrHxKBwGmOOSpphjYkcswF449j6EjsjbinTZbp8PJtjsVK1XFJStdzXdy/jaddAep2FGY+wyFAQ==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "linux" ], "engines": { - "node": ">=12" + "node": ">=18" + } + }, + "node_modules/@esbuild/netbsd-arm64": { + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/netbsd-arm64/-/netbsd-arm64-0.25.11.tgz", + "integrity": "sha512-hr9Oxj1Fa4r04dNpWr3P8QKVVsjQhqrMSUzZzf+LZcYjZNqhA3IAfPQdEh1FLVUJSiu6sgAwp3OmwBfbFgG2Xg==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "netbsd" + ], + "engines": { + "node": ">=18" } }, "node_modules/@esbuild/netbsd-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/netbsd-x64/-/netbsd-x64-0.18.20.tgz", - "integrity": "sha512-iO1c++VP6xUBUmltHZoMtCUdPlnPGdBom6IrO4gyKPFFVBKioIImVooR5I83nTew5UOYrk3gIJhbZh8X44y06A==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/netbsd-x64/-/netbsd-x64-0.25.11.tgz", + "integrity": "sha512-u7tKA+qbzBydyj0vgpu+5h5AeudxOAGncb8N6C9Kh1N4n7wU1Xw1JDApsRjpShRpXRQlJLb9wY28ELpwdPcZ7A==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "netbsd" ], "engines": { - "node": ">=12" + "node": ">=18" + } + }, + "node_modules/@esbuild/openbsd-arm64": { + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/openbsd-arm64/-/openbsd-arm64-0.25.11.tgz", + "integrity": "sha512-Qq6YHhayieor3DxFOoYM1q0q1uMFYb7cSpLD2qzDSvK1NAvqFi8Xgivv0cFC6J+hWVw2teCYltyy9/m/14ryHg==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "openbsd" + ], + "engines": { + "node": ">=18" } }, "node_modules/@esbuild/openbsd-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/openbsd-x64/-/openbsd-x64-0.18.20.tgz", - "integrity": "sha512-e5e4YSsuQfX4cxcygw/UCPIEP6wbIL+se3sxPdCiMbFLBWu0eiZOJ7WoD+ptCLrmjZBK1Wk7I6D/I3NglUGOxg==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/openbsd-x64/-/openbsd-x64-0.25.11.tgz", + "integrity": "sha512-CN+7c++kkbrckTOz5hrehxWN7uIhFFlmS/hqziSFVWpAzpWrQoAG4chH+nN3Be+Kzv/uuo7zhX716x3Sn2Jduw==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "openbsd" ], "engines": { - "node": ">=12" + "node": ">=18" + } + }, + "node_modules/@esbuild/openharmony-arm64": { + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/openharmony-arm64/-/openharmony-arm64-0.25.11.tgz", + "integrity": "sha512-rOREuNIQgaiR+9QuNkbkxubbp8MSO9rONmwP5nKncnWJ9v5jQ4JxFnLu4zDSRPf3x4u+2VN4pM4RdyIzDty/wQ==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "openharmony" + ], + "engines": { + "node": ">=18" } }, "node_modules/@esbuild/sunos-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/sunos-x64/-/sunos-x64-0.18.20.tgz", - "integrity": "sha512-kDbFRFp0YpTQVVrqUd5FTYmWo45zGaXe0X8E1G/LKFC0v8x0vWrhOWSLITcCn63lmZIxfOMXtCfti/RxN/0wnQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/sunos-x64/-/sunos-x64-0.25.11.tgz", + "integrity": "sha512-nq2xdYaWxyg9DcIyXkZhcYulC6pQ2FuCgem3LI92IwMgIZ69KHeY8T4Y88pcwoLIjbed8n36CyKoYRDygNSGhA==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "sunos" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/win32-arm64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/win32-arm64/-/win32-arm64-0.18.20.tgz", - "integrity": "sha512-ddYFR6ItYgoaq4v4JmQQaAI5s7npztfV4Ag6NrhiaW0RrnOXqBkgwZLofVTlq1daVTQNhtI5oieTvkRPfZrePg==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/win32-arm64/-/win32-arm64-0.25.11.tgz", + "integrity": "sha512-3XxECOWJq1qMZ3MN8srCJ/QfoLpL+VaxD/WfNRm1O3B4+AZ/BnLVgFbUV3eiRYDMXetciH16dwPbbHqwe1uU0Q==", "cpu": [ "arm64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "win32" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/win32-ia32": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/win32-ia32/-/win32-ia32-0.18.20.tgz", - "integrity": "sha512-Wv7QBi3ID/rROT08SABTS7eV4hX26sVduqDOTe1MvGMjNd3EjOz4b7zeexIR62GTIEKrfJXKL9LFxTYgkyeu7g==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/win32-ia32/-/win32-ia32-0.25.11.tgz", + "integrity": "sha512-3ukss6gb9XZ8TlRyJlgLn17ecsK4NSQTmdIXRASVsiS2sQ6zPPZklNJT5GR5tE/MUarymmy8kCEf5xPCNCqVOA==", "cpu": [ "ia32" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "win32" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@esbuild/win32-x64": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/@esbuild/win32-x64/-/win32-x64-0.18.20.tgz", - "integrity": "sha512-kTdfRcSiDfQca/y9QIkng02avJ+NCaQvrMejlsB3RRv5sE9rRoeBPISaZpKxHELzRxZyLvNts1P27W3wV+8geQ==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/@esbuild/win32-x64/-/win32-x64-0.25.11.tgz", + "integrity": "sha512-D7Hpz6A2L4hzsRpPaCYkQnGOotdUpDzSGRIv9I+1ITdHROSFUWW95ZPZWQmGka1Fg7W3zFJowyn9WGwMJ0+KPA==", "cpu": [ "x64" ], "dev": true, + "license": "MIT", "optional": true, "os": [ "win32" ], "engines": { - "node": ">=12" + "node": ">=18" } }, "node_modules/@loaders.gl/core": { "version": "4.2.2", "resolved": "https://registry.npmjs.org/@loaders.gl/core/-/core-4.2.2.tgz", "integrity": "sha512-d3YElSsqL29MaiOwzGB97v994SPotbTvJnooCqoQsYGoYYrECdIetv1/zlfDsh5UB2Wl/NaUMJrzyOqlLmDz5Q==", + "peer": true, "dependencies": { "@loaders.gl/loader-utils": "4.2.2", "@loaders.gl/schema": "4.2.2", @@ -453,12 +545,327 @@ "resolved": "https://registry.npmjs.org/@probe.gl/stats/-/stats-4.0.9.tgz", "integrity": "sha512-Q9Xt/sJUQaMsbjRKjOscv2t7wXIymTrOEJ4a3da4FTCn7bkKvcdxdyFAQySCrtPxE+YZ5I5lXpWPgv9BwmpE1g==" }, + "node_modules/@rollup/rollup-android-arm-eabi": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-android-arm-eabi/-/rollup-android-arm-eabi-4.52.5.tgz", + "integrity": "sha512-8c1vW4ocv3UOMp9K+gToY5zL2XiiVw3k7f1ksf4yO1FlDFQ1C2u72iACFnSOceJFsWskc2WZNqeRhFRPzv+wtQ==", + "cpu": [ + "arm" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "android" + ] + }, + "node_modules/@rollup/rollup-android-arm64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-android-arm64/-/rollup-android-arm64-4.52.5.tgz", + "integrity": "sha512-mQGfsIEFcu21mvqkEKKu2dYmtuSZOBMmAl5CFlPGLY94Vlcm+zWApK7F/eocsNzp8tKmbeBP8yXyAbx0XHsFNA==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "android" + ] + }, + "node_modules/@rollup/rollup-darwin-arm64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-darwin-arm64/-/rollup-darwin-arm64-4.52.5.tgz", + "integrity": "sha512-takF3CR71mCAGA+v794QUZ0b6ZSrgJkArC+gUiG6LB6TQty9T0Mqh3m2ImRBOxS2IeYBo4lKWIieSvnEk2OQWA==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "darwin" + ] + }, + "node_modules/@rollup/rollup-darwin-x64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-darwin-x64/-/rollup-darwin-x64-4.52.5.tgz", + "integrity": "sha512-W901Pla8Ya95WpxDn//VF9K9u2JbocwV/v75TE0YIHNTbhqUTv9w4VuQ9MaWlNOkkEfFwkdNhXgcLqPSmHy0fA==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "darwin" + ] + }, + "node_modules/@rollup/rollup-freebsd-arm64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-freebsd-arm64/-/rollup-freebsd-arm64-4.52.5.tgz", + "integrity": "sha512-QofO7i7JycsYOWxe0GFqhLmF6l1TqBswJMvICnRUjqCx8b47MTo46W8AoeQwiokAx3zVryVnxtBMcGcnX12LvA==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "freebsd" + ] + }, + "node_modules/@rollup/rollup-freebsd-x64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-freebsd-x64/-/rollup-freebsd-x64-4.52.5.tgz", + "integrity": "sha512-jr21b/99ew8ujZubPo9skbrItHEIE50WdV86cdSoRkKtmWa+DDr6fu2c/xyRT0F/WazZpam6kk7IHBerSL7LDQ==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "freebsd" + ] + }, + "node_modules/@rollup/rollup-linux-arm-gnueabihf": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm-gnueabihf/-/rollup-linux-arm-gnueabihf-4.52.5.tgz", + "integrity": "sha512-PsNAbcyv9CcecAUagQefwX8fQn9LQ4nZkpDboBOttmyffnInRy8R8dSg6hxxl2Re5QhHBf6FYIDhIj5v982ATQ==", + "cpu": [ + "arm" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-arm-musleabihf": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm-musleabihf/-/rollup-linux-arm-musleabihf-4.52.5.tgz", + "integrity": "sha512-Fw4tysRutyQc/wwkmcyoqFtJhh0u31K+Q6jYjeicsGJJ7bbEq8LwPWV/w0cnzOqR2m694/Af6hpFayLJZkG2VQ==", + "cpu": [ + "arm" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-arm64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm64-gnu/-/rollup-linux-arm64-gnu-4.52.5.tgz", + "integrity": "sha512-a+3wVnAYdQClOTlyapKmyI6BLPAFYs0JM8HRpgYZQO02rMR09ZcV9LbQB+NL6sljzG38869YqThrRnfPMCDtZg==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-arm64-musl": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-arm64-musl/-/rollup-linux-arm64-musl-4.52.5.tgz", + "integrity": "sha512-AvttBOMwO9Pcuuf7m9PkC1PUIKsfaAJ4AYhy944qeTJgQOqJYJ9oVl2nYgY7Rk0mkbsuOpCAYSs6wLYB2Xiw0Q==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-loong64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-loong64-gnu/-/rollup-linux-loong64-gnu-4.52.5.tgz", + "integrity": "sha512-DkDk8pmXQV2wVrF6oq5tONK6UHLz/XcEVow4JTTerdeV1uqPeHxwcg7aFsfnSm9L+OO8WJsWotKM2JJPMWrQtA==", + "cpu": [ + "loong64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-ppc64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-ppc64-gnu/-/rollup-linux-ppc64-gnu-4.52.5.tgz", + "integrity": "sha512-W/b9ZN/U9+hPQVvlGwjzi+Wy4xdoH2I8EjaCkMvzpI7wJUs8sWJ03Rq96jRnHkSrcHTpQe8h5Tg3ZzUPGauvAw==", + "cpu": [ + "ppc64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-riscv64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-riscv64-gnu/-/rollup-linux-riscv64-gnu-4.52.5.tgz", + "integrity": "sha512-sjQLr9BW7R/ZiXnQiWPkErNfLMkkWIoCz7YMn27HldKsADEKa5WYdobaa1hmN6slu9oWQbB6/jFpJ+P2IkVrmw==", + "cpu": [ + "riscv64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-riscv64-musl": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-riscv64-musl/-/rollup-linux-riscv64-musl-4.52.5.tgz", + "integrity": "sha512-hq3jU/kGyjXWTvAh2awn8oHroCbrPm8JqM7RUpKjalIRWWXE01CQOf/tUNWNHjmbMHg/hmNCwc/Pz3k1T/j/Lg==", + "cpu": [ + "riscv64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-s390x-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-s390x-gnu/-/rollup-linux-s390x-gnu-4.52.5.tgz", + "integrity": "sha512-gn8kHOrku8D4NGHMK1Y7NA7INQTRdVOntt1OCYypZPRt6skGbddska44K8iocdpxHTMMNui5oH4elPH4QOLrFQ==", + "cpu": [ + "s390x" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-x64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-x64-gnu/-/rollup-linux-x64-gnu-4.52.5.tgz", + "integrity": "sha512-hXGLYpdhiNElzN770+H2nlx+jRog8TyynpTVzdlc6bndktjKWyZyiCsuDAlpd+j+W+WNqfcyAWz9HxxIGfZm1Q==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-linux-x64-musl": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-linux-x64-musl/-/rollup-linux-x64-musl-4.52.5.tgz", + "integrity": "sha512-arCGIcuNKjBoKAXD+y7XomR9gY6Mw7HnFBv5Rw7wQRvwYLR7gBAgV7Mb2QTyjXfTveBNFAtPt46/36vV9STLNg==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "linux" + ] + }, + "node_modules/@rollup/rollup-openharmony-arm64": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-openharmony-arm64/-/rollup-openharmony-arm64-4.52.5.tgz", + "integrity": "sha512-QoFqB6+/9Rly/RiPjaomPLmR/13cgkIGfA40LHly9zcH1S0bN2HVFYk3a1eAyHQyjs3ZJYlXvIGtcCs5tko9Cw==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "openharmony" + ] + }, + "node_modules/@rollup/rollup-win32-arm64-msvc": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-arm64-msvc/-/rollup-win32-arm64-msvc-4.52.5.tgz", + "integrity": "sha512-w0cDWVR6MlTstla1cIfOGyl8+qb93FlAVutcor14Gf5Md5ap5ySfQ7R9S/NjNaMLSFdUnKGEasmVnu3lCMqB7w==", + "cpu": [ + "arm64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "win32" + ] + }, + "node_modules/@rollup/rollup-win32-ia32-msvc": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-ia32-msvc/-/rollup-win32-ia32-msvc-4.52.5.tgz", + "integrity": "sha512-Aufdpzp7DpOTULJCuvzqcItSGDH73pF3ko/f+ckJhxQyHtp67rHw3HMNxoIdDMUITJESNE6a8uh4Lo4SLouOUg==", + "cpu": [ + "ia32" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "win32" + ] + }, + "node_modules/@rollup/rollup-win32-x64-gnu": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-x64-gnu/-/rollup-win32-x64-gnu-4.52.5.tgz", + "integrity": "sha512-UGBUGPFp1vkj6p8wCRraqNhqwX/4kNQPS57BCFc8wYh0g94iVIW33wJtQAx3G7vrjjNtRaxiMUylM0ktp/TRSQ==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "win32" + ] + }, + "node_modules/@rollup/rollup-win32-x64-msvc": { + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/@rollup/rollup-win32-x64-msvc/-/rollup-win32-x64-msvc-4.52.5.tgz", + "integrity": "sha512-TAcgQh2sSkykPRWLrdyy2AiceMckNf5loITqXxFI5VuQjS5tSuw3WlwdN8qv8vzjLAUTvYaH/mVjSFpbkFbpTg==", + "cpu": [ + "x64" + ], + "dev": true, + "license": "MIT", + "optional": true, + "os": [ + "win32" + ] + }, "node_modules/@tweakpane/core": { "version": "1.1.9", "resolved": "https://registry.npmjs.org/@tweakpane/core/-/core-1.1.9.tgz", "integrity": "sha512-9tq+KAhaqPiOgsFyLPAz1IMXkVfhRqxGzAgy1ps3As6o3W7XjnU7sev6OlD/Z+Pzw8uZVMukkSHf2e0uCU6u0A==", "dev": true }, + "node_modules/@types/estree": { + "version": "1.0.8", + "resolved": "https://registry.npmjs.org/@types/estree/-/estree-1.0.8.tgz", + "integrity": "sha512-dWHzHa2WqEXI/O1E9OjrocMTKJl2mSrEolh1Iomrv6U+JuNwaHXsXx9bLu5gG7BUWFIN0skIQJQ/L1rIex4X6w==", + "dev": true, + "license": "MIT" + }, "node_modules/@types/geojson": { "version": "7946.0.14", "resolved": "https://registry.npmjs.org/@types/geojson/-/geojson-7946.0.14.tgz", @@ -471,40 +878,63 @@ "dev": true }, "node_modules/esbuild": { - "version": "0.18.20", - "resolved": "https://registry.npmjs.org/esbuild/-/esbuild-0.18.20.tgz", - "integrity": "sha512-ceqxoedUrcayh7Y7ZX6NdbbDzGROiyVBgC4PriJThBKSVPWnnFHZAkfI1lJT8QFkOwH4qOS2SJkS4wvpGl8BpA==", + "version": "0.25.11", + "resolved": "https://registry.npmjs.org/esbuild/-/esbuild-0.25.11.tgz", + "integrity": "sha512-KohQwyzrKTQmhXDW1PjCv3Tyspn9n5GcY2RTDqeORIdIJY8yKIF7sTSopFmn/wpMPW4rdPXI0UE5LJLuq3bx0Q==", "dev": true, "hasInstallScript": true, + "license": "MIT", "bin": { "esbuild": "bin/esbuild" }, "engines": { - "node": ">=12" + "node": ">=18" }, "optionalDependencies": { - "@esbuild/android-arm": "0.18.20", - "@esbuild/android-arm64": "0.18.20", - "@esbuild/android-x64": "0.18.20", - "@esbuild/darwin-arm64": "0.18.20", - "@esbuild/darwin-x64": "0.18.20", - "@esbuild/freebsd-arm64": "0.18.20", - "@esbuild/freebsd-x64": "0.18.20", - "@esbuild/linux-arm": "0.18.20", - "@esbuild/linux-arm64": "0.18.20", - "@esbuild/linux-ia32": "0.18.20", - "@esbuild/linux-loong64": "0.18.20", - "@esbuild/linux-mips64el": "0.18.20", - "@esbuild/linux-ppc64": "0.18.20", - "@esbuild/linux-riscv64": "0.18.20", - "@esbuild/linux-s390x": "0.18.20", - "@esbuild/linux-x64": "0.18.20", - "@esbuild/netbsd-x64": "0.18.20", - "@esbuild/openbsd-x64": "0.18.20", - "@esbuild/sunos-x64": "0.18.20", - "@esbuild/win32-arm64": "0.18.20", - "@esbuild/win32-ia32": "0.18.20", - "@esbuild/win32-x64": "0.18.20" + "@esbuild/aix-ppc64": "0.25.11", + "@esbuild/android-arm": "0.25.11", + "@esbuild/android-arm64": "0.25.11", + "@esbuild/android-x64": "0.25.11", + "@esbuild/darwin-arm64": "0.25.11", + "@esbuild/darwin-x64": "0.25.11", + "@esbuild/freebsd-arm64": "0.25.11", + "@esbuild/freebsd-x64": "0.25.11", + "@esbuild/linux-arm": "0.25.11", + "@esbuild/linux-arm64": "0.25.11", + "@esbuild/linux-ia32": "0.25.11", + "@esbuild/linux-loong64": "0.25.11", + "@esbuild/linux-mips64el": "0.25.11", + "@esbuild/linux-ppc64": "0.25.11", + "@esbuild/linux-riscv64": "0.25.11", + "@esbuild/linux-s390x": "0.25.11", + "@esbuild/linux-x64": "0.25.11", + "@esbuild/netbsd-arm64": "0.25.11", + "@esbuild/netbsd-x64": "0.25.11", + "@esbuild/openbsd-arm64": "0.25.11", + "@esbuild/openbsd-x64": "0.25.11", + "@esbuild/openharmony-arm64": "0.25.11", + "@esbuild/sunos-x64": "0.25.11", + "@esbuild/win32-arm64": "0.25.11", + "@esbuild/win32-ia32": "0.25.11", + "@esbuild/win32-x64": "0.25.11" + } + }, + "node_modules/fdir": { + "version": "6.5.0", + "resolved": "https://registry.npmjs.org/fdir/-/fdir-6.5.0.tgz", + "integrity": "sha512-tIbYtZbucOs0BRGqPJkshJUYdL+SDH7dVM8gjy+ERp3WAUjLEFJE+02kanyHtwjWOnwrKYBiwAmM0p4kLJAnXg==", + "dev": true, + "license": "MIT", + "engines": { + "node": ">=12.0.0" + }, + "peerDependencies": { + "picomatch": "^3 || ^4" + }, + "peerDependenciesMeta": { + "picomatch": { + "optional": true + } } }, "node_modules/fsevents": { @@ -513,6 +943,7 @@ "integrity": "sha512-5xoDfX+fL7faATnagmWPpbFtwh/R77WmMMqqHGS65C3vvB0YHrgF+B1YmZ3441tMj5n63k0212XNoJwzlhffQw==", "dev": true, "hasInstallScript": true, + "license": "MIT", "optional": true, "os": [ "darwin" @@ -522,9 +953,9 @@ } }, "node_modules/nanoid": { - "version": "3.3.7", - "resolved": "https://registry.npmjs.org/nanoid/-/nanoid-3.3.7.tgz", - "integrity": "sha512-eSRppjcPIatRIMC1U6UngP8XFcz8MQWGQdt1MTBQ7NaAmvXDfvNxbvWV3x2y6CdEUciCSsDHDQZbhYaB8QEo2g==", + "version": "3.3.11", + "resolved": "https://registry.npmjs.org/nanoid/-/nanoid-3.3.11.tgz", + "integrity": "sha512-N8SpfPUnUp1bK+PMYW8qSWdl9U+wwNWI4QKxOYDy9JAro3WMX7p2OeVRF9v+347pnakNevPmiHhNmZ2HbFA76w==", "dev": true, "funding": [ { @@ -532,6 +963,7 @@ "url": "https://github.com/sponsors/ai" } ], + "license": "MIT", "bin": { "nanoid": "bin/nanoid.cjs" }, @@ -540,15 +972,30 @@ } }, "node_modules/picocolors": { - "version": "1.0.1", - "resolved": "https://registry.npmjs.org/picocolors/-/picocolors-1.0.1.tgz", - "integrity": "sha512-anP1Z8qwhkbmu7MFP5iTt+wQKXgwzf7zTyGlcdzabySa9vd0Xt392U0rVmz9poOaBj0uHJKyyo9/upk0HrEQew==", - "dev": true + "version": "1.1.1", + "resolved": "https://registry.npmjs.org/picocolors/-/picocolors-1.1.1.tgz", + "integrity": "sha512-xceH2snhtb5M9liqDsmEw56le376mTZkEX/jEb/RxNFyegNul7eNslCXP9FDj/Lcu0X8KEyMceP2ntpaHrDEVA==", + "dev": true, + "license": "ISC" + }, + "node_modules/picomatch": { + "version": "4.0.3", + "resolved": "https://registry.npmjs.org/picomatch/-/picomatch-4.0.3.tgz", + "integrity": "sha512-5gTmgEY/sqK6gFXLIsQNH19lWb4ebPDLA4SdLP7dsWkIXHWlG66oPuVvXSGFPppYZz8ZDZq0dYYrbHfBCVUb1Q==", + "dev": true, + "license": "MIT", + "peer": true, + "engines": { + "node": ">=12" + }, + "funding": { + "url": "https://github.com/sponsors/jonschlinkert" + } }, "node_modules/postcss": { - "version": "8.4.41", - "resolved": "https://registry.npmjs.org/postcss/-/postcss-8.4.41.tgz", - "integrity": "sha512-TesUflQ0WKZqAvg52PWL6kHgLKP6xB6heTOdoYM0Wt2UHyxNa4K25EZZMgKns3BH1RLVbZCREPpLY0rhnNoHVQ==", + "version": "8.5.6", + "resolved": "https://registry.npmjs.org/postcss/-/postcss-8.5.6.tgz", + "integrity": "sha512-3Ybi1tAuwAP9s0r1UQ2J4n5Y0G05bJkpUIO0/bI9MhwmD70S5aTWbXGBwxHrelT+XM1k6dM0pk+SwNkpTRN7Pg==", "dev": true, "funding": [ { @@ -564,44 +1011,90 @@ "url": "https://github.com/sponsors/ai" } ], + "license": "MIT", "dependencies": { - "nanoid": "^3.3.7", - "picocolors": "^1.0.1", - "source-map-js": "^1.2.0" + "nanoid": "^3.3.11", + "picocolors": "^1.1.1", + "source-map-js": "^1.2.1" }, "engines": { "node": "^10 || ^12 || >=14" } }, "node_modules/rollup": { - "version": "3.29.5", - "resolved": "https://registry.npmjs.org/rollup/-/rollup-3.29.5.tgz", - "integrity": "sha512-GVsDdsbJzzy4S/v3dqWPJ7EfvZJfCHiDqe80IyrF59LYuP+e6U1LJoUqeuqRbwAWoMNoXivMNeNAOf5E22VA1w==", + "version": "4.52.5", + "resolved": "https://registry.npmjs.org/rollup/-/rollup-4.52.5.tgz", + "integrity": "sha512-3GuObel8h7Kqdjt0gxkEzaifHTqLVW56Y/bjN7PSQtkKr0w3V/QYSdt6QWYtd7A1xUtYQigtdUfgj1RvWVtorw==", "dev": true, + "license": "MIT", + "dependencies": { + "@types/estree": "1.0.8" + }, "bin": { "rollup": "dist/bin/rollup" }, "engines": { - "node": ">=14.18.0", + "node": ">=18.0.0", "npm": ">=8.0.0" }, "optionalDependencies": { + "@rollup/rollup-android-arm-eabi": "4.52.5", + "@rollup/rollup-android-arm64": "4.52.5", + "@rollup/rollup-darwin-arm64": "4.52.5", + "@rollup/rollup-darwin-x64": "4.52.5", + "@rollup/rollup-freebsd-arm64": "4.52.5", + "@rollup/rollup-freebsd-x64": "4.52.5", + "@rollup/rollup-linux-arm-gnueabihf": "4.52.5", + "@rollup/rollup-linux-arm-musleabihf": "4.52.5", + "@rollup/rollup-linux-arm64-gnu": "4.52.5", + "@rollup/rollup-linux-arm64-musl": "4.52.5", + "@rollup/rollup-linux-loong64-gnu": "4.52.5", + "@rollup/rollup-linux-ppc64-gnu": "4.52.5", + "@rollup/rollup-linux-riscv64-gnu": "4.52.5", + "@rollup/rollup-linux-riscv64-musl": "4.52.5", + "@rollup/rollup-linux-s390x-gnu": "4.52.5", + "@rollup/rollup-linux-x64-gnu": "4.52.5", + "@rollup/rollup-linux-x64-musl": "4.52.5", + "@rollup/rollup-openharmony-arm64": "4.52.5", + "@rollup/rollup-win32-arm64-msvc": "4.52.5", + "@rollup/rollup-win32-ia32-msvc": "4.52.5", + "@rollup/rollup-win32-x64-gnu": "4.52.5", + "@rollup/rollup-win32-x64-msvc": "4.52.5", "fsevents": "~2.3.2" } }, "node_modules/source-map-js": { - "version": "1.2.0", - "resolved": "https://registry.npmjs.org/source-map-js/-/source-map-js-1.2.0.tgz", - "integrity": "sha512-itJW8lvSA0TXEphiRoawsCksnlf8SyvmFzIhltqAHluXd88pkCd+cXJVHTDwdCr0IzwptSm035IHQktUu1QUMg==", + "version": "1.2.1", + "resolved": "https://registry.npmjs.org/source-map-js/-/source-map-js-1.2.1.tgz", + "integrity": "sha512-UXWMKhLOwVKb728IUtQPXxfYU+usdybtUrK/8uGE8CQMvrhOpwvzDBwj0QhSL7MQc7vIsISBG8VQ8+IDQxpfQA==", "dev": true, + "license": "BSD-3-Clause", "engines": { "node": ">=0.10.0" } }, + "node_modules/tinyglobby": { + "version": "0.2.15", + "resolved": "https://registry.npmjs.org/tinyglobby/-/tinyglobby-0.2.15.tgz", + "integrity": "sha512-j2Zq4NyQYG5XMST4cbs02Ak8iJUdxRM0XI5QyxXuZOzKOINmWurp3smXu3y5wDcJrptwpSjgXHzIQxR0omXljQ==", + "dev": true, + "license": "MIT", + "dependencies": { + "fdir": "^6.5.0", + "picomatch": "^4.0.3" + }, + "engines": { + "node": ">=12.0.0" + }, + "funding": { + "url": "https://github.com/sponsors/SuperchupuDev" + } + }, "node_modules/tweakpane": { "version": "3.1.10", "resolved": "https://registry.npmjs.org/tweakpane/-/tweakpane-3.1.10.tgz", "integrity": "sha512-rqwnl/pUa7+inhI2E9ayGTqqP0EPOOn/wVvSWjZsRbZUItzNShny7pzwL3hVlaN4m9t/aZhsP0aFQ9U5VVR2VQ==", + "peer": true, "funding": { "url": "https://github.com/sponsors/cocopon" } @@ -629,40 +1122,51 @@ } }, "node_modules/vite": { - "version": "4.5.5", - "resolved": "https://registry.npmjs.org/vite/-/vite-4.5.5.tgz", - "integrity": "sha512-ifW3Lb2sMdX+WU91s3R0FyQlAyLxOzCSCP37ujw0+r5POeHPwe6udWVIElKQq8gk3t7b8rkmvqC6IHBpCff4GQ==", + "version": "7.1.12", + "resolved": "https://registry.npmjs.org/vite/-/vite-7.1.12.tgz", + "integrity": "sha512-ZWyE8YXEXqJrrSLvYgrRP7p62OziLW7xI5HYGWFzOvupfAlrLvURSzv/FyGyy0eidogEM3ujU+kUG1zuHgb6Ug==", "dev": true, + "license": "MIT", "dependencies": { - "esbuild": "^0.18.10", - "postcss": "^8.4.27", - "rollup": "^3.27.1" + "esbuild": "^0.25.0", + "fdir": "^6.5.0", + "picomatch": "^4.0.3", + "postcss": "^8.5.6", + "rollup": "^4.43.0", + "tinyglobby": "^0.2.15" }, "bin": { "vite": "bin/vite.js" }, "engines": { - "node": "^14.18.0 || >=16.0.0" + "node": "^20.19.0 || >=22.12.0" }, "funding": { "url": "https://github.com/vitejs/vite?sponsor=1" }, "optionalDependencies": { - "fsevents": "~2.3.2" + "fsevents": "~2.3.3" }, "peerDependencies": { - "@types/node": ">= 14", - "less": "*", + "@types/node": "^20.19.0 || >=22.12.0", + "jiti": ">=1.21.0", + "less": "^4.0.0", "lightningcss": "^1.21.0", - "sass": "*", - "stylus": "*", - "sugarss": "*", - "terser": "^5.4.0" + "sass": "^1.70.0", + "sass-embedded": "^1.70.0", + "stylus": ">=0.54.8", + "sugarss": "^5.0.0", + "terser": "^5.16.0", + "tsx": "^4.8.1", + "yaml": "^2.4.2" }, "peerDependenciesMeta": { "@types/node": { "optional": true }, + "jiti": { + "optional": true + }, "less": { "optional": true }, @@ -672,6 +1176,9 @@ "sass": { "optional": true }, + "sass-embedded": { + "optional": true + }, "stylus": { "optional": true }, @@ -680,6 +1187,12 @@ }, "terser": { "optional": true + }, + "tsx": { + "optional": true + }, + "yaml": { + "optional": true } } }, diff --git a/package.json b/package.json index 8a751f0..2a4bb7b 100644 --- a/package.json +++ b/package.json @@ -9,10 +9,10 @@ }, "devDependencies": { "@tweakpane/core": "^1.1.7", - "tweakpane-plugin-file-import": "^0.2.0", "@webgpu/types": "^0.1.31", + "tweakpane-plugin-file-import": "^0.2.0", "typescript": "^5.0.4", - "vite": "^4.3.1", + "vite": "^7.1.12", "vite-raw-plugin": "^1.0.1" }, "dependencies": { diff --git a/src/renderers/gaussian-renderer-tile.ts b/src/renderers/gaussian-renderer-tile.ts new file mode 100644 index 0000000..83efbd5 --- /dev/null +++ b/src/renderers/gaussian-renderer-tile.ts @@ -0,0 +1,261 @@ +// Gaussian splatting renderer with tile-based depth sorting +// Complete separate pipeline for per-tile depth ordering + +import { PointCloud } from '../utils/load'; +import commonWSGL from '../shaders/common.wgsl' +import preprocessTileWGSL from '../shaders/preprocess-tile.wgsl'; +import renderWGSL from '../shaders/gaussian.wgsl'; +import tileDepthSortWGSL from '../sort/tile-depth-sort.wgsl'; +import { get_sorter, c_histogram_block_rows, C } from '../sort/sort'; +import { create_tile_depth_sorter } from '../sort/tile-depth-sort'; +import { Renderer } from './renderer'; + +export interface GaussianTileAwareRenderer extends Renderer { + set_gaussian_multiplier: (v: number) => void; + set_tile_size: (size: number) => void; +} + +const createBuffer = ( + device: GPUDevice, + label: string, + size: number, + usage: GPUBufferUsageFlags, + data?: ArrayBuffer | ArrayBufferView +) => { + const buffer = device.createBuffer({ label, size, usage }); + if (data) device.queue.writeBuffer(buffer, 0, data as ArrayBuffer); + return buffer; +}; + +export default function get_renderer( + pc: PointCloud, + device: GPUDevice, + presentation_format: GPUTextureFormat, + camera_buffer: GPUBuffer, + canvas: HTMLCanvasElement, +): GaussianTileAwareRenderer { + + const sorter = get_sorter(pc.num_points, device); + + const nulling_data = new Uint32Array([0]); + + const draw_indirect_buffer = createBuffer( + device, + 'draw indirect buffer tile', + 4 * Uint32Array.BYTES_PER_ELEMENT, + GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, + new Uint32Array([6, pc.num_points, 0, 0]) + ); + + const SPLAT_SIZE = 24; + const splat_buffer = createBuffer( + device, + 'splats buffer tile', + pc.num_points * SPLAT_SIZE, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + ); + + const settings_buffer = createBuffer( + device, + 'render settings buffer tile', + 4 * Uint32Array.BYTES_PER_ELEMENT, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Float32Array([1.0, pc.sh_deg, 0.0, 0.0]) + ); + + let tile_size = 32; + const tile_config_buffer = createBuffer( + device, + 'tile config', + 16, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Uint32Array([ + tile_size, + Math.ceil(canvas.width / tile_size), + Math.ceil(canvas.height / tile_size), + Math.ceil(canvas.width / tile_size) * Math.ceil(canvas.height / tile_size) + ]) + ); + + const tile_touched_buffer = createBuffer( + device, + 'tile touched', + pc.num_points * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const max_tile_pairs = Math.min(pc.num_points * 8, 16 * 1024 * 1024); + const tile_sorter = create_tile_depth_sorter( + pc.num_points, + max_tile_pairs, + canvas.width, + canvas.height, + tile_size, + device, + commonWSGL + tileDepthSortWGSL + ); + + const preprocess_pipeline = device.createComputePipeline({ + label: 'preprocess tile', + layout: 'auto', + compute: { + module: device.createShaderModule({ code: commonWSGL + preprocessTileWGSL }), + entryPoint: 'preprocess', + constants: { + workgroupSize: C.histogram_wg_size, + sortKeyPerThread: c_histogram_block_rows, + }, + }, + }); + + const sort_bind_group = device.createBindGroup({ + label: 'sort tile', + layout: preprocess_pipeline.getBindGroupLayout(2), + entries: [ + { binding: 0, resource: { buffer: sorter.sort_info_buffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_depths_buffer } }, + { binding: 2, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + { binding: 3, resource: { buffer: sorter.sort_dispatch_indirect_buffer } }, + ], + }); + + const tile_bind_group = device.createBindGroup({ + label: 'tile config bind', + layout: preprocess_pipeline.getBindGroupLayout(3), + entries: [ + { binding: 0, resource: { buffer: tile_config_buffer } }, + { binding: 1, resource: { buffer: tile_touched_buffer } }, + ], + }); + + const preprocess_bind_group0 = device.createBindGroup({ + label: 'preprocess 0 tile', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } }, + { binding: 1, resource: { buffer: pc.gaussian_3d_buffer } }, + { binding: 2, resource: { buffer: pc.sh_buffer } }, + ], + }); + + const preprocess_bind_group1 = device.createBindGroup({ + label: 'preprocess 1 tile', + layout: preprocess_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: settings_buffer } }, + ], + }); + + const render_shader = device.createShaderModule({ code: commonWSGL + renderWGSL }); + const render_pipeline = device.createRenderPipeline({ + label: 'gaussian render tile', + layout: 'auto', + vertex: { + module: render_shader, + entryPoint: 'vs_main', + }, + fragment: { + module: render_shader, + entryPoint: 'fs_main', + targets: [{ + format: presentation_format, + blend: { + color: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + alpha: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + }, + writeMask: GPUColorWrite.ALL, + }], + }, + primitive: { + topology: 'triangle-list', + }, + }); + + const render_bind_group = device.createBindGroup({ + label: 'gaussian splats tile', + layout: render_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + ], + }); + + function record_preprocess(encoder: GPUCommandEncoder) { + device.queue.writeBuffer(sorter.sort_info_buffer, 0, nulling_data); + device.queue.writeBuffer(sorter.sort_dispatch_indirect_buffer, 0, nulling_data); + + const pass = encoder.beginComputePass({ label: 'preprocess tile' }); + pass.setPipeline(preprocess_pipeline); + pass.setBindGroup(0, preprocess_bind_group0); + pass.setBindGroup(1, preprocess_bind_group1); + pass.setBindGroup(2, sort_bind_group); + pass.setBindGroup(3, tile_bind_group); + const wg = Math.ceil(pc.num_points / C.histogram_wg_size); + pass.dispatchWorkgroups(wg); + pass.end(); + } + + function record_tile_sort(encoder: GPUCommandEncoder) { + tile_sorter.key_generation_pass( + encoder, + camera_buffer, + splat_buffer, + sorter.ping_pong[0].sort_depths_buffer, + tile_sorter.tile_offsets_buffer, + tile_touched_buffer + ); + sorter.sort(encoder); + tile_sorter.identify_ranges_pass(encoder); + } + + function record_render(encoder: GPUCommandEncoder, texture_view: GPUTextureView) { + const pass = encoder.beginRenderPass({ + label: 'gaussian render tile', + colorAttachments: [ + { + view: texture_view, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(render_pipeline); + pass.setBindGroup(0, render_bind_group); + pass.drawIndirect(draw_indirect_buffer, 0); + pass.end(); + } + + + return { + frame: (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + record_preprocess(encoder); + record_tile_sort(encoder); + encoder.copyBufferToBuffer(sorter.sort_info_buffer, 0, draw_indirect_buffer, 4, 4); + record_render(encoder, texture_view); + }, + camera_buffer, + set_gaussian_multiplier(v: number) { + device.queue.writeBuffer(settings_buffer, 0, new Float32Array([v, pc.sh_deg, 0.0, 0.0])); + }, + set_tile_size(size: number) { + tile_size = size; + const num_tiles_x = Math.ceil(canvas.width / size); + const num_tiles_y = Math.ceil(canvas.height / size); + device.queue.writeBuffer( + tile_config_buffer, + 0, + new Uint32Array([size, num_tiles_x, num_tiles_y, num_tiles_x * num_tiles_y]) + ); + console.log(`[TileRenderer] Set tile_size=${size}`); + }, + }; +} diff --git a/src/renderers/gaussian-renderer.ts b/src/renderers/gaussian-renderer.ts index 1684523..eadd676 100644 --- a/src/renderers/gaussian-renderer.ts +++ b/src/renderers/gaussian-renderer.ts @@ -1,11 +1,12 @@ import { PointCloud } from '../utils/load'; +import commonWSGL from '../shaders/common.wgsl' import preprocessWGSL from '../shaders/preprocess.wgsl'; import renderWGSL from '../shaders/gaussian.wgsl'; import { get_sorter,c_histogram_block_rows,C } from '../sort/sort'; import { Renderer } from './renderer'; export interface GaussianRenderer extends Renderer { - + set_gaussian_multiplier: (v: number) => void; } // Utility to create GPU buffers @@ -17,7 +18,7 @@ const createBuffer = ( data?: ArrayBuffer | ArrayBufferView ) => { const buffer = device.createBuffer({ label, size, usage }); - if (data) device.queue.writeBuffer(buffer, 0, data); + if (data) device.queue.writeBuffer(buffer, 0, data as ArrayBuffer); return buffer; }; @@ -36,6 +37,30 @@ export default function get_renderer( const nulling_data = new Uint32Array([0]); + const draw_indirect_buffer = createBuffer( + device, + 'draw indirect buffer', + 4 * Uint32Array.BYTES_PER_ELEMENT, + GPUBufferUsage.INDIRECT | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC, + new Uint32Array([6, pc.num_points, 0, 0]) + ); + + const SPLAT_SIZE = 24; + const splat_buffer = createBuffer( + device, + 'splats buffer', + pc.num_points * SPLAT_SIZE, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + ); + + const settings_buffer = createBuffer( + device, + 'render settings buffer', + 4 * Uint32Array.BYTES_PER_ELEMENT, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Float32Array([1.0, pc.sh_deg, 0.0, 0.0]) + ); + // =============================================== // Create Compute Pipeline and Bind Groups // =============================================== @@ -43,7 +68,7 @@ export default function get_renderer( label: 'preprocess', layout: 'auto', compute: { - module: device.createShaderModule({ code: preprocessWGSL }), + module: device.createShaderModule({ code: commonWSGL + preprocessWGSL }), entryPoint: 'preprocess', constants: { workgroupSize: C.histogram_wg_size, @@ -67,11 +92,101 @@ export default function get_renderer( // =============================================== // Create Render Pipeline and Bind Groups // =============================================== - + const preprocess_bind_group0 = device.createBindGroup({ + label: 'preprocess 0', + layout: preprocess_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } }, + { binding: 1, resource: { buffer: pc.gaussian_3d_buffer } }, + { binding: 2, resource: { buffer: pc.sh_buffer } }, + ], + }); + + const preprocess_bind_group1 = device.createBindGroup({ + label: 'preprocess 1', + layout: preprocess_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: settings_buffer } }, + ], + }); + + const render_shader = device.createShaderModule({ code: commonWSGL + renderWGSL }); + const render_pipeline = device.createRenderPipeline({ + label: 'gaussian render', + layout: 'auto', + vertex: { + module: render_shader, + entryPoint: 'vs_main', + }, + fragment: { + module: render_shader, + entryPoint: 'fs_main', + targets: [{ + format: presentation_format, + blend: { + color: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + alpha: { + srcFactor: 'one', + dstFactor: 'one-minus-src-alpha', + operation: 'add', + }, + }, + writeMask: GPUColorWrite.ALL, + }], + }, + primitive: { + topology: 'triangle-list', + }, + }); + + const render_bind_group = device.createBindGroup({ + label: 'gaussian splats', + layout: render_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: splat_buffer } }, + { binding: 1, resource: { buffer: sorter.ping_pong[0].sort_indices_buffer } }, + ], + }); + // =============================================== // Command Encoder Functions // =============================================== + function record_preprocess(encoder: GPUCommandEncoder) { + device.queue.writeBuffer(sorter.sort_info_buffer, 0, nulling_data); + device.queue.writeBuffer(sorter.sort_dispatch_indirect_buffer, 0, nulling_data); + + const pass = encoder.beginComputePass({ label: 'preprocess' }); + pass.setPipeline(preprocess_pipeline); + pass.setBindGroup(0, preprocess_bind_group0); + pass.setBindGroup(1, preprocess_bind_group1); + pass.setBindGroup(2, sort_bind_group); + const wg = Math.ceil(pc.num_points / C.histogram_wg_size); + pass.dispatchWorkgroups(wg); + pass.end(); + } + + function record_render(encoder: GPUCommandEncoder, texture_view: GPUTextureView) { + const pass = encoder.beginRenderPass({ + label: 'gaussian render', + colorAttachments: [ + { + view: texture_view, + loadOp: 'clear', + storeOp: 'store', + }, + ], + }); + pass.setPipeline(render_pipeline); + pass.setBindGroup(0, render_bind_group); + pass.drawIndirect(draw_indirect_buffer, 0); + pass.end(); + } // =============================================== @@ -79,8 +194,14 @@ export default function get_renderer( // =============================================== return { frame: (encoder: GPUCommandEncoder, texture_view: GPUTextureView) => { + record_preprocess(encoder); sorter.sort(encoder); + encoder.copyBufferToBuffer(sorter.sort_info_buffer, 0, draw_indirect_buffer, 4, 4); + record_render(encoder, texture_view); }, camera_buffer, + set_gaussian_multiplier(v: number) { + device.queue.writeBuffer(settings_buffer, 0, new Float32Array([v, pc.sh_deg, 0.0, 0.0])); + }, }; } diff --git a/src/renderers/renderer.ts b/src/renderers/renderer.ts index ffdf9ba..d8df0af 100644 --- a/src/renderers/renderer.ts +++ b/src/renderers/renderer.ts @@ -2,6 +2,7 @@ import { load } from '../utils/load'; import { Pane } from 'tweakpane'; import * as TweakpaneFileImportPlugin from 'tweakpane-plugin-file-import'; import { default as get_renderer_gaussian, GaussianRenderer } from './gaussian-renderer'; +import { default as get_renderer_gaussian_tile, GaussianTileAwareRenderer } from './gaussian-renderer-tile'; import { default as get_renderer_pointcloud } from './point-cloud-renderer'; import { Camera, load_camera_presets} from '../camera/camera'; import { CameraControl } from '../camera/camera-control'; @@ -19,8 +20,9 @@ export default async function init( ) { let ply_file_loaded = false; let cam_file_loaded = false; - let renderers: { pointcloud?: Renderer, gaussian?: Renderer } = {}; + let renderers: { pointcloud?: Renderer, gaussian?: Renderer, gaussian_tile?: Renderer } = {}; let gaussian_renderer: GaussianRenderer | undefined; + let gaussian_tile_renderer: GaussianTileAwareRenderer | undefined; let pointcloud_renderer: Renderer | undefined; let renderer: Renderer | undefined; let cameras; @@ -48,6 +50,7 @@ export default async function init( const params = { fps: 0.0, gaussian_multiplier: 1, + tile_size: 16, renderer: 'pointcloud', ply_file: '', cam_file: '', @@ -68,6 +71,7 @@ export default async function init( options: { pointcloud: 'pointcloud', gaussian: 'gaussian', + gaussian_tile: 'gaussian_tile', } }).on('change', (e) => { renderer = renderers[e.value]; @@ -86,9 +90,11 @@ export default async function init( const pc = await load(uploadedFile, device); pointcloud_renderer = get_renderer_pointcloud(pc, device, presentation_format, camera.uniform_buffer); gaussian_renderer = get_renderer_gaussian(pc, device, presentation_format, camera.uniform_buffer); + gaussian_tile_renderer = get_renderer_gaussian_tile(pc, device, presentation_format, camera.uniform_buffer, canvas); renderers = { pointcloud: pointcloud_renderer, gaussian: gaussian_renderer, + gaussian_tile: gaussian_tile_renderer, }; renderer = renderers[params.renderer]; ply_file_loaded = true; @@ -121,7 +127,17 @@ export default async function init( 'gaussian_multiplier', {min: 0, max: 1.5} ).on('change', (e) => { - //TODO: Bind constants to the gaussian renderer. + gaussian_renderer?.set_gaussian_multiplier(e.value); + gaussian_tile_renderer?.set_gaussian_multiplier(e.value); + }); + } + { + pane.addInput( + params, + 'tile_size', + {min: 4, max: 128, step: 4} + ).on('change', (e) => { + gaussian_tile_renderer?.set_tile_size(e.value); }); } diff --git a/src/shaders/common.wgsl b/src/shaders/common.wgsl new file mode 100644 index 0000000..e47db7e --- /dev/null +++ b/src/shaders/common.wgsl @@ -0,0 +1,27 @@ +struct CameraUniforms { + view: mat4x4, + view_inv: mat4x4, + proj: mat4x4, + proj_inv: mat4x4, + viewport: vec2, + focal: vec2 +}; + +struct RenderSettings { + gaussian_scaling: f32, + sh_deg: f32, +} + +struct Gaussian { + pos_opacity: array, + rot: array, + scale: array +}; + +struct Splat { + //TODO: store information for 2D splat rendering + center: u32, + size: u32, + rgba: array, + conic: array, +}; diff --git a/src/shaders/gaussian.wgsl b/src/shaders/gaussian.wgsl index 759226d..4a7898f 100644 --- a/src/shaders/gaussian.wgsl +++ b/src/shaders/gaussian.wgsl @@ -1,22 +1,71 @@ struct VertexOutput { @builtin(position) position: vec4, //TODO: information passed from vertex shader to fragment shader + @location(0) rgba: vec4, + @location(1) opacity: f32, + @location(2) center: vec2, + @location(3) conic: vec3, }; -struct Splat { - //TODO: information defined in preprocess compute shader -}; +@group(0) @binding(0) +var splats : array; +@group(0) @binding(1) +var sort_indices : array; + +const NORMALS: array,6> = array,6>( + vec2(-1.0f, -1.0f), + vec2( 1.0f, -1.0f), + vec2( 1.0f, 1.0f), + vec2(-1.0f, -1.0f), + vec2( 1.0f, 1.0f), + vec2(-1.0f, 1.0f), +); @vertex fn vs_main( + @builtin(vertex_index) vertex_index: u32, + @builtin(instance_index) instance_index: u32 ) -> VertexOutput { //TODO: reconstruct 2D quad based on information from splat, pass var out: VertexOutput; - out.position = vec4(1. ,1. , 0., 1.); + let index = sort_indices[instance_index]; + let splat = splats[index]; + + let center = unpack2x16float(splat.center); + let size = unpack2x16float(splat.size); + let rg = unpack2x16float(splat.rgba[0]); + let ba = unpack2x16float(splat.rgba[1]); + let conic_xy = unpack2x16float(splat.conic[0]); + let conic_zw = unpack2x16float(splat.conic[1]); + + let conic = vec3(conic_xy, conic_zw.x); + let radius = conic_zw.y; + let offset = NORMALS[vertex_index % 6u] * radius / size; + + out.position = vec4(center + offset, 0.0f, 1.0f); + out.rgba = vec4(rg, ba.x, 1.0f); + out.opacity = ba.y; + out.center = vec2(center.x + 1.0f, -center.y + 1.0f) * size - 1.0f; + out.conic = conic; return out; } @fragment fn fs_main(in: VertexOutput) -> @location(0) vec4 { - return vec4(1.); + let x = in.conic.x; + let y = in.conic.y; + let z = in.conic.z; + + let offset = in.position.xy - in.center; + let rx = offset.x; + let ry = offset.y; + + let prod = x * rx * rx + z * ry * ry + 2.0f * y * rx * ry; + + if (prod <= 0.0f) { + discard; + } + + let opacity = clamp(in.opacity * exp(-0.5f * prod), 0.0f, 0.99f); + return in.rgba * opacity; } \ No newline at end of file diff --git a/src/shaders/point_cloud.wgsl b/src/shaders/point_cloud.wgsl index 01dded1..eb34b7b 100644 --- a/src/shaders/point_cloud.wgsl +++ b/src/shaders/point_cloud.wgsl @@ -35,7 +35,8 @@ fn vs_main( let pos = vec4(a.x, a.y, b.x, 1.); // TODO: MVP calculations - out.position = pos; + let view_pos = camera.view * pos; + out.position = camera.proj * view_pos; return out; } diff --git a/src/shaders/preprocess-tile.wgsl b/src/shaders/preprocess-tile.wgsl new file mode 100644 index 0000000..4a26e8b --- /dev/null +++ b/src/shaders/preprocess-tile.wgsl @@ -0,0 +1,248 @@ +// Preprocessing for tile-based depth sorting +// Computes splat properties and tile offsets for per-tile rendering + +const SH_C0: f32 = 0.28209479177387814; +const SH_C1 = 0.4886025119029199; +const SH_C2 = array( + 1.0925484305920792, + -1.0925484305920792, + 0.31539156525252005, + -1.0925484305920792, + 0.5462742152960396 +); +const SH_C3 = array( + -0.5900435899266435, + 2.890611442640554, + -0.4570457994644658, + 0.3731763325901154, + -0.4570457994644658, + 1.445305721320277, + -0.5900435899266435 +); + +override workgroupSize: u32; +override sortKeyPerThread: u32; + +struct DispatchIndirect { + dispatch_x: atomic, + dispatch_y: u32, + dispatch_z: u32, +} + +struct SortInfos { + keys_size: atomic, + padded_size: u32, + passes: u32, + even_pass: u32, + odd_pass: u32, +} + +struct TileConfig { + tile_size: u32, + num_tiles_x: u32, + num_tiles_y: u32, + total_tiles: u32, +}; + +@group(0) @binding(0) +var camera: CameraUniforms; +@group(0) @binding(1) +var gaussians : array; +@group(0) @binding(2) +var sh_buf : array; + +@group(1) @binding(0) +var splats : array; +@group(1) @binding(1) +var settings: RenderSettings; + +@group(2) @binding(0) +var sort_infos: SortInfos; +@group(2) @binding(1) +var sort_depths : array; +@group(2) @binding(2) +var sort_indices : array; +@group(2) @binding(3) +var sort_dispatch: DispatchIndirect; + +@group(3) @binding(0) +var tile_config: TileConfig; +@group(3) @binding(1) +var tile_touched: array; + +fn sh_coef(splat_idx: u32, c_idx: u32) -> vec3 { + let resid = c_idx % 2u; + let ind = splat_idx * 24u + (c_idx >> 1u) * 3u + resid; + let rg = unpack2x16float(sh_buf[ind]); + let ba = unpack2x16float(sh_buf[ind + 1u]); + if (resid == 0u) { + return vec3(rg, ba.x); + } + return vec3(rg.y, ba); +} + +fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { + var result = SH_C0 * sh_coef(v_idx, 0u); + + if sh_deg > 0u { + let x = dir.x; + let y = dir.y; + let z = dir.z; + result += - SH_C1 * y * sh_coef(v_idx, 1u) + SH_C1 * z * sh_coef(v_idx, 2u) - SH_C1 * x * sh_coef(v_idx, 3u); + + if sh_deg > 1u { + let xx = dir.x * dir.x; + let yy = dir.y * dir.y; + let zz = dir.z * dir.z; + let xy = dir.x * dir.y; + let yz = dir.y * dir.z; + let xz = dir.x * dir.z; + + result += SH_C2[0] * xy * sh_coef(v_idx, 4u) + SH_C2[1] * yz * sh_coef(v_idx, 5u) + SH_C2[2] * (2.0 * zz - xx - yy) * sh_coef(v_idx, 6u) + SH_C2[3] * xz * sh_coef(v_idx, 7u) + SH_C2[4] * (xx - yy) * sh_coef(v_idx, 8u); + + if sh_deg > 2u { + result += SH_C3[0] * y * (3.0 * xx - yy) * sh_coef(v_idx, 9u) + SH_C3[1] * xy * z * sh_coef(v_idx, 10u) + SH_C3[2] * y * (4.0 * zz - xx - yy) * sh_coef(v_idx, 11u) + SH_C3[3] * z * (2.0 * zz - 3.0 * xx - 3.0 * yy) * sh_coef(v_idx, 12u) + SH_C3[4] * x * (4.0 * zz - xx - yy) * sh_coef(v_idx, 13u) + SH_C3[5] * z * (xx - yy) * sh_coef(v_idx, 14u) + SH_C3[6] * x * (xx - 3.0 * yy) * sh_coef(v_idx, 15u); + } + } + } + result += 0.5; + return max(vec3(0.), result); +} + +fn quaternion_to_rotation_matrix(q: vec4) -> mat3x3 { + let r = q.x; + let x = q.y; + let y = q.z; + let z = q.w; + return mat3x3( + 1.f - 2.f * (y * y + z * z), 2.f * (x * y - r * z), 2.f * (x * z + r * y), + 2.f * (x * y + r * z), 1.f - 2.f * (x * x + z * z), 2.f * (y * z - r * x), + 2.f * (x * z - r * y), 2.f * (y * z + r * x), 1.f - 2.f * (x * x + y * y), + ); +} + +fn view_position_to_jacobian_matrix(pos_view: vec4, focal: vec2) -> mat3x3 { + let x_view = pos_view.x; + let y_view = pos_view.y; + let z_view = pos_view.z; + let fx = focal.x; + let fy = focal.y; + return mat3x3( + fx / z_view, 0.0f, -fx * x_view / (z_view * z_view), + 0.0f, fy / z_view, -fy * y_view / (z_view * z_view), + 0.0f, 0.0f, 0.0f, + ); +} + +fn cov3d_to_cov2d( + mat_cov3d: mat3x3, + mat_jacobian: mat3x3, + view_matrix: mat4x4 +) -> vec3 { + let mat_view = mat3x3(view_matrix[0].xyz, view_matrix[1].xyz, view_matrix[2].xyz); + let mat_t = transpose(mat_view) * mat_jacobian; + let mat_cov = transpose(mat_t) * mat_cov3d * mat_t; + return vec3(mat_cov[0][0] + 0.3f, mat_cov[0][1], mat_cov[1][1] + 0.3f); +} + +fn get_tile_rect(center: vec2, radius: i32) -> vec4 { + let min_x = max(0i, i32(center.x) - radius) / i32(tile_config.tile_size); + let min_y = max(0i, i32(center.y) - radius) / i32(tile_config.tile_size); + let max_x = (i32(center.x) + radius + i32(tile_config.tile_size) - 1) / i32(tile_config.tile_size); + let max_y = (i32(center.y) + radius + i32(tile_config.tile_size) - 1) / i32(tile_config.tile_size); + + return vec4( + u32(clamp(min_x, 0i, i32(tile_config.num_tiles_x))), + u32(clamp(min_y, 0i, i32(tile_config.num_tiles_y))), + u32(clamp(max_x, 0i, i32(tile_config.num_tiles_x))), + u32(clamp(max_y, 0i, i32(tile_config.num_tiles_y))) + ); +} + +@compute @workgroup_size(workgroupSize,1,1) +fn preprocess(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgroups) wgs: vec3) { + let idx = gid.x; + if (idx >= arrayLength(&gaussians)) { + return; + } + + let gaussian = gaussians[idx]; + let xy = unpack2x16float(gaussian.pos_opacity[0]); + let zw = unpack2x16float(gaussian.pos_opacity[1]); + let pos_world = vec4(xy.x, xy.y, zw.x, 1.0f); + let pos_view = camera.view * pos_world; + + if (pos_view.z < 1e-3f) { + return; + } + + let pos_clip = camera.proj * pos_view; + let pos_ndc = pos_clip.xyz / (pos_clip.w + 1e-6f); + + if (abs(pos_ndc.x) > 1.2f || abs(pos_ndc.y) > 1.2f) { + return; + } + + let s01 = unpack2x16float(gaussian.scale[0]); + let s23 = unpack2x16float(gaussian.scale[1]); + let scale = exp(vec3(s01.x, s01.y, s23.x)) * settings.gaussian_scaling; + let mat_scale = mat3x3( + scale.x, 0.0f, 0.0f, + 0.0f, scale.y, 0.0f, + 0.0f, 0.0f, scale.z, + ); + + let quaternion = normalize( + vec4( + unpack2x16float(gaussian.rot[0]), + unpack2x16float(gaussian.rot[1]), + ) + ); + let mat_rotation = quaternion_to_rotation_matrix(quaternion); + + let mat_jacobian = view_position_to_jacobian_matrix(pos_view, camera.focal); + + let mat_m = mat_scale * mat_rotation; + let mat_cov3d = transpose(mat_m) * mat_m; + + let vec_cov2d = cov3d_to_cov2d(mat_cov3d, mat_jacobian, camera.view); + + let a = vec_cov2d.x; + let b = vec_cov2d.y; + let c = vec_cov2d.z; + let det = (a * c - b * b); + if (det <= 0.0f) { + return; + } + let mid = 0.5f * (a + c); + let lambda1 = mid + sqrt(max(0.1f, mid * mid - det)); + let lambda2 = mid - sqrt(max(0.1f, mid * mid - det)); + let radius = i32(ceil(3.0f * sqrt(max(lambda1, lambda2)))); + + let det_inv = 1.0f / det; + let conic = vec3(c * det_inv, -b * det_inv, a * det_inv); + + let rgb = computeColorFromSH(normalize(pos_view.xyz), idx, u32(settings.sh_deg)); + let opacity = 1.0f / (1.0f + exp(-zw.y)); + + let write_idx = atomicAdd(&sort_infos.keys_size, 1u); + + sort_indices[write_idx] = write_idx; + sort_depths[write_idx] = 0xFFFFFFFFu ^ bitcast(-pos_view.z); + + splats[write_idx].center = pack2x16float(pos_ndc.xy); + splats[write_idx].size = pack2x16float(0.5f * camera.viewport); + splats[write_idx].rgba[0] = pack2x16float(rgb.xy); + splats[write_idx].rgba[1] = pack2x16float(vec2(rgb.z, opacity)); + splats[write_idx].conic[0] = pack2x16float(conic.xy); + splats[write_idx].conic[1] = pack2x16float(vec2(conic.z, f32(radius))); + + let keys_per_dispatch = workgroupSize * sortKeyPerThread; + if ((write_idx % keys_per_dispatch) == 0u) { + _ = atomicAdd(&sort_dispatch.dispatch_x, 1u); + } + + let rect = get_tile_rect(pos_ndc.xy * camera.viewport, radius); + let num_tiles = (rect.z - rect.x) * (rect.w - rect.y); + tile_touched[idx] = num_tiles; +} diff --git a/src/shaders/preprocess.wgsl b/src/shaders/preprocess.wgsl index bbc63f5..514e96d 100644 --- a/src/shaders/preprocess.wgsl +++ b/src/shaders/preprocess.wgsl @@ -35,30 +35,6 @@ struct SortInfos { odd_pass: u32, } -struct CameraUniforms { - view: mat4x4, - view_inv: mat4x4, - proj: mat4x4, - proj_inv: mat4x4, - viewport: vec2, - focal: vec2 -}; - -struct RenderSettings { - gaussian_scaling: f32, - sh_deg: f32, -} - -struct Gaussian { - pos_opacity: array, - rot: array, - scale: array -}; - -struct Splat { - //TODO: store information for 2D splat rendering -}; - //TODO: bind your data here @group(2) @binding(0) var sort_infos: SortInfos; @@ -69,10 +45,29 @@ var sort_indices : array; @group(2) @binding(3) var sort_dispatch: DispatchIndirect; +@group(0) @binding(0) +var camera: CameraUniforms; +@group(0) @binding(1) +var gaussians : array; +@group(0) @binding(2) +var sh_buf : array; + +@group(1) @binding(0) +var splats : array; +@group(1) @binding(1) +var settings: RenderSettings; + /// reads the ith sh coef from the storage buffer fn sh_coef(splat_idx: u32, c_idx: u32) -> vec3 { //TODO: access your binded sh_coeff, see load.ts for how it is stored - return vec3(0.0); + let resid = c_idx % 2u; + let ind = splat_idx * 24u + (c_idx >> 1u) * 3u + resid; + let rg = unpack2x16float(sh_buf[ind]); + let ba = unpack2x16float(sh_buf[ind + 1u]); + if (resid == 0u) { + return vec3(rg, ba.x); + } + return vec3(rg.y, ba); } // spherical harmonics evaluation with Condon–Shortley phase @@ -108,11 +103,131 @@ fn computeColorFromSH(dir: vec3, v_idx: u32, sh_deg: u32) -> vec3 { return max(vec3(0.), result); } +fn quaternion_to_rotation_matrix(q: vec4) -> mat3x3 { + let r = q.x; + let x = q.y; + let y = q.z; + let z = q.w; + return mat3x3( + 1.f - 2.f * (y * y + z * z), 2.f * (x * y - r * z), 2.f * (x * z + r * y), + 2.f * (x * y + r * z), 1.f - 2.f * (x * x + z * z), 2.f * (y * z - r * x), + 2.f * (x * z - r * y), 2.f * (y * z + r * x), 1.f - 2.f * (x * x + y * y), + ); +} + +fn view_position_to_jacobian_matrix(pos_view: vec4, focal: vec2) -> mat3x3 { + let x_view = pos_view.x; + let y_view = pos_view.y; + let z_view = pos_view.z; + let fx = focal.x; + let fy = focal.y; + return mat3x3( + fx / z_view, 0.0f, -fx * x_view / (z_view * z_view), + 0.0f, fy / z_view, -fy * y_view / (z_view * z_view), + 0.0f, 0.0f, 0.0f, + ); +} + +fn cov3d_to_cov2d( + mat_cov3d: mat3x3, + mat_jacobian: mat3x3, + view_matrix: mat4x4 +) -> vec3 { + let mat_view = mat3x3(view_matrix[0].xyz, view_matrix[1].xyz, view_matrix[2].xyz); + let mat_t = transpose(mat_view) * mat_jacobian; + let mat_cov = transpose(mat_t) * mat_cov3d * mat_t; + return vec3(mat_cov[0][0] + 0.3f, mat_cov[0][1], mat_cov[1][1] + 0.3f); +} + @compute @workgroup_size(workgroupSize,1,1) fn preprocess(@builtin(global_invocation_id) gid: vec3, @builtin(num_workgroups) wgs: vec3) { let idx = gid.x; - //TODO: set up pipeline as described in instruction + if (idx >= arrayLength(&gaussians)) { + return; + } + + let gaussian = gaussians[idx]; + let xy = unpack2x16float(gaussian.pos_opacity[0]); + let zw = unpack2x16float(gaussian.pos_opacity[1]); + let pos_world = vec4(xy.x, xy.y, zw.x, 1.0f); + let pos_view = camera.view * pos_world; + + // z culling + if (pos_view.z < 1e-3f) { + return; + } + + let pos_clip = camera.proj * pos_view; + let pos_ndc = pos_clip.xyz / (pos_clip.w + 1e-6f); + + // xy culling + if (abs(pos_ndc.x) > 1.2f || abs(pos_ndc.y) > 1.2f) { + return; + } + + // scaling + let s01 = unpack2x16float(gaussian.scale[0]); + let s23 = unpack2x16float(gaussian.scale[1]); + let scale = exp(vec3(s01.x, s01.y, s23.x)) * settings.gaussian_scaling; + let mat_scale = mat3x3( + scale.x, 0.0f, 0.0f, + 0.0f, scale.y, 0.0f, + 0.0f, 0.0f, scale.z, + ); + + // rotation + let quaternion = normalize( + vec4( + unpack2x16float(gaussian.rot[0]), + unpack2x16float(gaussian.rot[1]), + ) + ); + let mat_rotation = quaternion_to_rotation_matrix(quaternion); + + // jacobian + let mat_jacobian = view_position_to_jacobian_matrix(pos_view, camera.focal); + + // cov3d + let mat_m = mat_scale * mat_rotation; + let mat_cov3d = transpose(mat_m) * mat_m; + + // cov2d + let vec_cov2d = cov3d_to_cov2d(mat_cov3d, mat_jacobian, camera.view); + + // radius + let a = vec_cov2d.x; + let b = vec_cov2d.y; + let c = vec_cov2d.z; + let det = (a * c - b * b); + if (det <= 0.0f) { + return; + } + let mid = 0.5f * (a + c); + let lambda1 = mid + sqrt(max(0.1f, mid * mid - det)); + let lambda2 = mid - sqrt(max(0.1f, mid * mid - det)); + let radius = ceil(3.0f * sqrt(max(lambda1, lambda2))); + + let det_inv = 1.0f / det; + let conic = vec3(c * det_inv, -b * det_inv, a * det_inv); + + let rgb = computeColorFromSH(normalize(pos_view.xyz), idx, u32(settings.sh_deg)); + let opacity = 1.0f / (1.0f + exp(-zw.y)); + + let write_idx = atomicAdd(&sort_infos.keys_size, 1u); + + sort_indices[write_idx] = write_idx; + sort_depths[write_idx] = 0xFFFFFFFFu ^ bitcast(-pos_view.z); + + splats[write_idx].center = pack2x16float(pos_ndc.xy); + splats[write_idx].size = pack2x16float(0.5f * camera.viewport); + splats[write_idx].rgba[0] = pack2x16float(rgb.xy); + splats[write_idx].rgba[1] = pack2x16float(vec2(rgb.z, opacity)); + splats[write_idx].conic[0] = pack2x16float(conic.xy); + splats[write_idx].conic[1] = pack2x16float(vec2(conic.z, radius)); let keys_per_dispatch = workgroupSize * sortKeyPerThread; // increment DispatchIndirect.dispatchx each time you reach limit for one dispatch of keys + if ((write_idx % keys_per_dispatch) == 0u) { + _ = atomicAdd(&sort_dispatch.dispatch_x, 1u); + } } \ No newline at end of file diff --git a/src/sort/tile-depth-sort.ts b/src/sort/tile-depth-sort.ts new file mode 100644 index 0000000..02cbe5d --- /dev/null +++ b/src/sort/tile-depth-sort.ts @@ -0,0 +1,235 @@ +// Tile-based depth sorting pipeline +// Generates tile-sorted indices for per-tile Gaussian rendering + +export interface TileDepthSorter { + tile_size: number; + num_tiles_x: number; + num_tiles_y: number; + max_tile_pairs: number; + + tile_config_buffer: GPUBuffer; + tile_offsets_buffer: GPUBuffer; + tile_keys_hi_buffer: GPUBuffer; + tile_keys_lo_buffer: GPUBuffer; + tile_values_buffer: GPUBuffer; + tile_ranges_buffer: GPUBuffer; + tile_write_counter_buffer: GPUBuffer; + + key_generation_pipeline: GPUComputePipeline; + key_generation_bind_group_config: GPUBindGroup; + key_generation_bind_group_output: GPUBindGroup; + + identify_ranges_pipeline: GPUComputePipeline; + + key_generation_pass: (encoder: GPUCommandEncoder, camera: GPUBuffer, splat_data: GPUBuffer, sort_depths: GPUBuffer, tile_offsets: GPUBuffer, tile_touched: GPUBuffer) => void; + identify_ranges_pass: (encoder: GPUCommandEncoder) => void; +} + +const createBuffer = ( + device: GPUDevice, + label: string, + size: number, + usage: GPUBufferUsageFlags, + data?: ArrayBuffer | ArrayBufferView +): GPUBuffer => { + const buffer = device.createBuffer({ label, size, usage }); + if (data) device.queue.writeBuffer(buffer, 0, data as ArrayBuffer); + return buffer; +}; + +export function create_tile_depth_sorter( + num_gaussians: number, + max_tile_pairs: number, + canvas_width: number, + canvas_height: number, + tile_size: number, + device: GPUDevice, + shader_code: string +): TileDepthSorter { + const num_tiles_x = Math.ceil(canvas_width / tile_size); + const num_tiles_y = Math.ceil(canvas_height / tile_size); + + console.log(`[TileDepthSort] Initialized: ${num_gaussians} gaussians, tile_size=${tile_size}, grid=${num_tiles_x}x${num_tiles_y}, max_pairs=${max_tile_pairs}`); + + const tile_config_buffer = createBuffer( + device, + 'tile config', + 16, + GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST, + new Uint32Array([tile_size, num_tiles_x, num_tiles_y, num_tiles_x * num_tiles_y]) + ); + + const tile_offsets_buffer = createBuffer( + device, + 'tile offsets', + num_gaussians * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const tile_keys_hi_buffer = createBuffer( + device, + 'tile keys hi', + max_tile_pairs * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const tile_keys_lo_buffer = createBuffer( + device, + 'tile keys lo', + max_tile_pairs * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const tile_values_buffer = createBuffer( + device, + 'tile values', + max_tile_pairs * 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST + ); + + const tile_ranges_buffer = createBuffer( + device, + 'tile ranges', + num_tiles_x * num_tiles_y * 8, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + new Uint32Array(num_tiles_x * num_tiles_y * 2) + ); + + const tile_write_counter_buffer = createBuffer( + device, + 'tile write counter', + 4, + GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + new Uint32Array([0]) + ); + + const shader_module = device.createShaderModule({ code: shader_code }); + + const key_generation_pipeline = device.createComputePipeline({ + label: 'tile key generation', + layout: 'auto', + compute: { + module: shader_module, + entryPoint: 'key_generation', + }, + }); + + const identify_ranges_pipeline = device.createComputePipeline({ + label: 'tile identify ranges', + layout: 'auto', + compute: { + module: shader_module, + entryPoint: 'identify_ranges', + }, + }); + + const key_generation_bind_group_config = device.createBindGroup({ + label: 'tile key generation config', + layout: key_generation_pipeline.getBindGroupLayout(1), + entries: [ + { binding: 0, resource: { buffer: tile_config_buffer } }, + ], + }); + + const key_generation_bind_group_output = device.createBindGroup({ + label: 'tile key generation output', + layout: key_generation_pipeline.getBindGroupLayout(2), + entries: [ + { binding: 0, resource: { buffer: tile_keys_hi_buffer } }, + { binding: 1, resource: { buffer: tile_keys_lo_buffer } }, + { binding: 2, resource: { buffer: tile_values_buffer } }, + ], + }); + + function key_generation_pass( + encoder: GPUCommandEncoder, + camera_buffer: GPUBuffer, + splat_data: GPUBuffer, + sort_depths_buffer: GPUBuffer, + tile_offsets_buffer: GPUBuffer, + tile_touched_buffer: GPUBuffer + ) { + device.queue.writeBuffer(tile_write_counter_buffer, 0, new Uint32Array([0])); + + const bg_splats = device.createBindGroup({ + label: 'tile key generation splats (frame)', + layout: key_generation_pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: camera_buffer } }, + { binding: 1, resource: { buffer: splat_data } }, + { binding: 2, resource: { buffer: sort_depths_buffer } }, + ], + }); + + const bg_counters = device.createBindGroup({ + label: 'tile key generation counters (frame)', + layout: key_generation_pipeline.getBindGroupLayout(3), + entries: [ + { binding: 1, resource: { buffer: tile_write_counter_buffer } }, + ], + }); + + const pass = encoder.beginComputePass({ label: 'tile key generation' }); + pass.setPipeline(key_generation_pipeline); + pass.setBindGroup(0, bg_splats); + pass.setBindGroup(1, key_generation_bind_group_config); + pass.setBindGroup(2, key_generation_bind_group_output); + pass.setBindGroup(3, bg_counters); + const wg = Math.ceil(num_gaussians / 256); + pass.dispatchWorkgroups(wg); + pass.end(); + } + + function identify_ranges_pass(encoder: GPUCommandEncoder) { + device.queue.writeBuffer(tile_ranges_buffer, 0, new Uint32Array(num_tiles_x * num_tiles_y * 2)); + + const bg_keys = device.createBindGroup({ + label: 'tile identify ranges keys (frame)', + layout: identify_ranges_pipeline.getBindGroupLayout(2), + entries: [ + { binding: 0, resource: { buffer: tile_keys_hi_buffer } }, + { binding: 3, resource: { buffer: tile_ranges_buffer } }, + ], + }); + + const bg_counters = device.createBindGroup({ + label: 'tile identify ranges counters (frame)', + layout: identify_ranges_pipeline.getBindGroupLayout(3), + entries: [ + { binding: 1, resource: { buffer: tile_write_counter_buffer } }, + ], + }); + + const pass = encoder.beginComputePass({ label: 'tile identify ranges' }); + pass.setPipeline(identify_ranges_pipeline); + pass.setBindGroup(2, bg_keys); + pass.setBindGroup(3, bg_counters); + const wg = Math.min(Math.ceil(max_tile_pairs / 256), 65535); + pass.dispatchWorkgroups(wg); + pass.end(); + } + + return { + tile_size, + num_tiles_x, + num_tiles_y, + max_tile_pairs, + + tile_config_buffer, + tile_offsets_buffer, + tile_keys_hi_buffer, + tile_keys_lo_buffer, + tile_values_buffer, + tile_ranges_buffer, + tile_write_counter_buffer, + + key_generation_pipeline, + key_generation_bind_group_config, + key_generation_bind_group_output, + + identify_ranges_pipeline, + + key_generation_pass, + identify_ranges_pass, + }; +} diff --git a/src/sort/tile-depth-sort.wgsl b/src/sort/tile-depth-sort.wgsl new file mode 100644 index 0000000..c266adf --- /dev/null +++ b/src/sort/tile-depth-sort.wgsl @@ -0,0 +1,121 @@ +struct TileConfig { + tile_size: u32, + num_tiles_x: u32, + num_tiles_y: u32, + total_tiles: u32, +}; + +@group(0) @binding(0) +var camera: CameraUniforms; + +@group(0) @binding(1) +var splats: array; + +@group(0) @binding(2) +var sort_depths: array; + +@group(1) @binding(0) +var tile_config: TileConfig; + +@group(2) @binding(0) +var tile_keys_hi: array; + +@group(2) @binding(1) +var tile_keys_lo: array; + +@group(2) @binding(2) +var tile_values: array; + +@group(2) @binding(3) +var tile_ranges: array>; + +@group(3) @binding(0) +var tile_touched: array; + +@group(3) @binding(1) +var tile_write_counter: atomic; + +fn get_rect_min_max(center_ndc: vec2, radius: i32, viewport: vec2) -> vec4 { + let center_screen = (center_ndc + 1.0) * viewport * 0.5; + + let min_x = max(0i, i32(center_screen.x) - radius) / i32(tile_config.tile_size); + let min_y = max(0i, i32(center_screen.y) - radius) / i32(tile_config.tile_size); + let max_x = (i32(center_screen.x) + radius + i32(tile_config.tile_size) - 1) / i32(tile_config.tile_size); + let max_y = (i32(center_screen.y) + radius + i32(tile_config.tile_size) - 1) / i32(tile_config.tile_size); + + return vec4( + u32(clamp(min_x, 0i, i32(tile_config.num_tiles_x))), + u32(clamp(min_y, 0i, i32(tile_config.num_tiles_y))), + u32(clamp(max_x, 0i, i32(tile_config.num_tiles_x))), + u32(clamp(max_y, 0i, i32(tile_config.num_tiles_y))) + ); +} + +@compute @workgroup_size(256, 1, 1) +fn key_generation(@builtin(global_invocation_id) gid: vec3) { + let idx = gid.x; + if idx >= arrayLength(&splats) { + return; + } + + let splat = splats[idx]; + let center_ndc = unpack2x16float(splat.center); + + let conic_z_and_radius = unpack2x16float(splat.conic[1]); + let radius = i32(conic_z_and_radius.y); + + if radius <= 0 { + return; + } + + let rect = get_rect_min_max(center_ndc, radius, camera.viewport); + + let num_tiles = (rect.z - rect.x) * (rect.w - rect.y); + if num_tiles == 0u { + return; + } + + let base_pos = atomicAdd(&tile_write_counter, num_tiles); + var write_pos = base_pos; + + for (var y = rect.y; y < rect.w; y = y + 1u) { + for (var x = rect.x; x < rect.z; x = x + 1u) { + if write_pos >= arrayLength(&tile_keys_hi) { + return; + } + + let tile_id = y * tile_config.num_tiles_x + x; + + tile_keys_hi[write_pos] = tile_id; + tile_keys_lo[write_pos] = sort_depths[idx]; + tile_values[write_pos] = idx; + write_pos = write_pos + 1u; + } + } +} + +@compute @workgroup_size(256, 1, 1) +fn identify_ranges(@builtin(global_invocation_id) gid: vec3) { + let idx = gid.x; + + let total_pairs = atomicLoad(&tile_write_counter); + if idx >= total_pairs { + return; + } + + let tile_id = tile_keys_hi[idx]; + + if idx == 0u { + tile_ranges[tile_id].x = 0u; + } else { + let prev_tile_id = tile_keys_hi[idx - 1u]; + if tile_id != prev_tile_id { + tile_ranges[prev_tile_id].y = idx; + tile_ranges[tile_id].x = idx; + } + } + + if idx == total_pairs - 1u { + tile_ranges[tile_id].y = total_pairs; + } +}