GTAO demo mod

CMakeLists.txt

@@ -559,7 +559,8 @@ include(cmake/ModSDK.cmake)
 
 if (DUSK_ENABLE_CODE_MODS AND CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
     add_custom_target(dusklight_mods) # Aggregate target for all in-tree mods
-    add_subdirectory(tools/mod_template)
+    add_subdirectory(mods/template_mod)
+    add_subdirectory(mods/ao_mod)
 endif ()
 
 if (APPLE)

cmake/GameABIConfig.cmake

@@ -25,3 +25,8 @@ set(_game_include_dirs
 add_library(dusklight_game_headers INTERFACE)
 target_include_directories(dusklight_game_headers INTERFACE ${_game_include_dirs})
 target_compile_definitions(dusklight_game_headers INTERFACE ${_game_compile_defs})
+if (TARGET dawn::dawncpp_headers)
+    target_link_libraries(dusklight_game_headers INTERFACE dawn::dawncpp_headers)
+elseif (TARGET dawn::webgpu_dawn)
+    target_link_libraries(dusklight_game_headers INTERFACE dawn::webgpu_dawn)
+endif ()

docs/modding.md

@@ -28,7 +28,7 @@ function, read and write data fields, and hook the vast majority of game functio
 
 ## Getting Started
 
-Fork the [mod template](../tools/mod_template/), a self-contained CMake project that uses the Dusklight mod SDK.
+Fork the [mod template](../mods/template_mod/), a self-contained CMake project that uses the Dusklight mod SDK.
 
 ```
 my_mod/

mods/ao_mod/CMakeLists.txt

@@ -0,0 +1,20 @@
+cmake_minimum_required(VERSION 3.25)
+project(ao_mod CXX)
+
+if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
+    set(DUSK_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../.." CACHE PATH "Path to dusk source root")
+    option(DUSK_MOD_USE_FULL_TREE "Use full build instead of the minimal mod SDK" OFF)
+    set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+    if (DUSK_MOD_USE_FULL_TREE)
+        add_subdirectory("${DUSK_DIR}" dusk EXCLUDE_FROM_ALL)
+    else ()
+        add_subdirectory("${DUSK_DIR}/sdk" dusk-sdk EXCLUDE_FROM_ALL)
+    endif ()
+endif ()
+
+add_mod(ao_mod
+    SOURCES  src/mod.cpp
+    MOD_JSON mod.json
+    RES_DIR  res
+    BUNDLE
+)

mods/ao_mod/mod.json

@@ -0,0 +1,7 @@
+{
+  "id": "dev.twilitrealm.ao_mod",
+  "name": "[Demo] Ambient Occlusion",
+  "version": "1.0.0",
+  "author": "Twilit Realm",
+  "description": "Ground-truth ambient occlusion (GTAO) computed from the scene depth buffer and composited over the game. Ported from Bevy Engine's SSAO and Intel XeGTAO."
+}

mods/ao_mod/res/composite.wgsl

@@ -0,0 +1,161 @@
+// Fullscreen composite: multiplies the denoised ambient-occlusion visibility over the scene.
+//
+// Debug views:
+//   1 = raw AO visibility as grayscale
+//   2 = view-space normals reconstructed from depth (keep in sync with gtao.wgsl)
+//   3 = the preprocessed depth input
+//   4 = depth staircase detector
+
+struct Uniforms {
+    projection: mat4x4f,
+    inverse_projection: mat4x4f,
+    size: vec2f,        // AO texture size in pixels (may be half the render size)
+    inv_size: vec2f,
+    depth_scale: vec2f,
+    effect_radius: f32,
+    intensity: f32,
+    slice_count: f32,
+    samples_per_slice_side: f32,
+    debug_view: u32,
+    _pad: f32,
+}
+
+@group(0) @binding(0) var ambient_occlusion: texture_2d<f32>;
+@group(0) @binding(1) var preprocessed_depth: texture_2d<f32>;
+@group(0) @binding(2) var scene_depth_raw: texture_2d<f32>;
+@group(0) @binding(3) var<uniform> uniforms: Uniforms;
+
+struct VertexOutput {
+    @builtin(position) position: vec4f,
+    @location(0) uv: vec2f,
+}
+
+@vertex
+fn vs_main(@builtin(vertex_index) index: u32) -> VertexOutput {
+    // Fullscreen triangle
+    var out: VertexOutput;
+    let uv = vec2f(f32((index << 1u) & 2u), f32(index & 2u));
+    out.position = vec4f(uv * vec2f(2.0, -2.0) + vec2f(-1.0, 1.0), 0.0, 1.0);
+    out.uv = uv;
+    return out;
+}
+
+// Manual bilinear sample (r32float is unfilterable without optional device features)
+fn sample_visibility(uv: vec2f) -> f32 {
+    let coordinates = uv * uniforms.size - 0.5;
+    let base = floor(coordinates);
+    let fraction = coordinates - base;
+    let max_coordinates = vec2i(uniforms.size) - 1i;
+    let p00 = clamp(vec2i(base), vec2i(0i), max_coordinates);
+    let p11 = clamp(vec2i(base) + 1i, vec2i(0i), max_coordinates);
+    let v00 = textureLoad(ambient_occlusion, vec2i(p00.x, p00.y), 0i).r;
+    let v10 = textureLoad(ambient_occlusion, vec2i(p11.x, p00.y), 0i).r;
+    let v01 = textureLoad(ambient_occlusion, vec2i(p00.x, p11.y), 0i).r;
+    let v11 = textureLoad(ambient_occlusion, vec2i(p11.x, p11.y), 0i).r;
+    let top = mix(v00, v10, fraction.x);
+    let bottom = mix(v01, v11, fraction.x);
+    return mix(top, bottom, fraction.y);
+}
+
+fn load_depth(pixel_coordinates: vec2<i32>) -> f32 {
+    let coordinates = clamp(pixel_coordinates, vec2<i32>(0i), vec2<i32>(uniforms.size) - 1i);
+    return textureLoad(preprocessed_depth, coordinates, 0i).r;
+}
+
+fn reconstruct_view_space_position(depth: f32, uv: vec2f) -> vec3f {
+    let clip_xy = vec2f(uv.x * 2.0 - 1.0, 1.0 - 2.0 * uv.y);
+    let t = uniforms.inverse_projection * vec4f(clip_xy, depth, 1.0);
+    return t.xyz / t.w;
+}
+
+fn view_position_at(pixel_coordinates: vec2<i32>) -> vec3f {
+    let depth = load_depth(pixel_coordinates);
+    let uv = (vec2f(pixel_coordinates) + 0.5) * uniforms.inv_size;
+    return reconstruct_view_space_position(depth, uv);
+}
+
+fn reconstruct_normal(pixel_coordinates: vec2<i32>, pixel_position: vec3f, depth_center: f32) -> vec3f {
+    let depth_left1 = load_depth(pixel_coordinates + vec2<i32>(-1i, 0i));
+    let depth_left2 = load_depth(pixel_coordinates + vec2<i32>(-2i, 0i));
+    let depth_right1 = load_depth(pixel_coordinates + vec2<i32>(1i, 0i));
+    let depth_right2 = load_depth(pixel_coordinates + vec2<i32>(2i, 0i));
+    let depth_top1 = load_depth(pixel_coordinates + vec2<i32>(0i, -1i));
+    let depth_top2 = load_depth(pixel_coordinates + vec2<i32>(0i, -2i));
+    let depth_bottom1 = load_depth(pixel_coordinates + vec2<i32>(0i, 1i));
+    let depth_bottom2 = load_depth(pixel_coordinates + vec2<i32>(0i, 2i));
+
+    let use_left = abs(2.0 * depth_left1 - depth_left2 - depth_center) <
+        abs(2.0 * depth_right1 - depth_right2 - depth_center);
+    let use_top = abs(2.0 * depth_top1 - depth_top2 - depth_center) <
+        abs(2.0 * depth_bottom1 - depth_bottom2 - depth_center);
+
+    var ddx: vec3f;
+    if use_left {
+        ddx = pixel_position - view_position_at(pixel_coordinates + vec2<i32>(-1i, 0i));
+    } else {
+        ddx = view_position_at(pixel_coordinates + vec2<i32>(1i, 0i)) - pixel_position;
+    }
+    var ddy: vec3f;
+    if use_top {
+        ddy = pixel_position - view_position_at(pixel_coordinates + vec2<i32>(0i, -1i));
+    } else {
+        ddy = view_position_at(pixel_coordinates + vec2<i32>(0i, 1i)) - pixel_position;
+    }
+
+    var normal = normalize(cross(ddy, ddx));
+    if dot(normal, pixel_position) > 0.0 {
+        normal = -normal;
+    }
+    return normal;
+}
+
+// Raw-snapshot variant of load_depth for the staircase view
+fn load_raw_depth(pixel_coordinates: vec2<i32>) -> f32 {
+    let size = vec2<i32>(textureDimensions(scene_depth_raw));
+    let coordinates = clamp(pixel_coordinates, vec2<i32>(0i), size - 1i);
+    return textureLoad(scene_depth_raw, coordinates, 0i).r;
+}
+
+@fragment
+fn fs_main(in: VertexOutput) -> @location(0) vec4f {
+    if uniforms.debug_view == 2u {
+        // Reconstructed view-space normals, [-1,1] -> RGB
+        let pixel = vec2<i32>(in.uv * uniforms.size);
+        let depth = load_depth(pixel);
+        let uv = (vec2f(pixel) + 0.5) * uniforms.inv_size;
+        let position = reconstruct_view_space_position(depth, uv);
+        let normal = reconstruct_normal(pixel, position, depth);
+        return vec4f(normal * 0.5 + 0.5, 1.0);
+    }
+    if uniforms.debug_view == 3u {
+        // Preprocessed depth as an exponential distance gradient (white = near, black = far)
+        let pixel = vec2<i32>(in.uv * uniforms.size);
+        let position = view_position_at(pixel);
+        let value = exp(-max(-position.z, 0.0) * 0.0003);
+        return vec4f(value, value, value, 1.0);
+    }
+    if uniforms.debug_view == 4u {
+        // Staircase detector on the raw snapshot depth
+        let size = vec2f(textureDimensions(scene_depth_raw));
+        let pixel = vec2<i32>(in.uv * size);
+        let d_center = load_raw_depth(pixel);
+        let d_left = load_raw_depth(pixel + vec2<i32>(-1i, 0i));
+        let d_right = load_raw_depth(pixel + vec2<i32>(1i, 0i));
+        let d_top = load_raw_depth(pixel + vec2<i32>(0i, -1i));
+        let d_bottom = load_raw_depth(pixel + vec2<i32>(0i, 1i));
+        let gradient_x = abs(d_right - d_left) * 0.5;
+        let curvature_x = abs(d_right - 2.0 * d_center + d_left);
+        let gradient_y = abs(d_bottom - d_top) * 0.5;
+        let curvature_y = abs(d_bottom - 2.0 * d_center + d_top);
+        let ratio_x = curvature_x / max(gradient_x, 1e-12);
+        let ratio_y = curvature_y / max(gradient_y, 1e-12);
+        return vec4f(saturate(ratio_x), saturate(ratio_y), 0.0, 1.0);
+    }
+
+    let visibility = sample_visibility(in.uv);
+    if uniforms.debug_view == 1u {
+        return vec4f(visibility, visibility, visibility, 1.0);
+    }
+    let value = mix(1.0, visibility, uniforms.intensity);
+    return vec4f(value, value, value, 1.0);
+}

mods/ao_mod/res/denoise.wgsl

@@ -0,0 +1,108 @@
+// 3x3 bilaterial filter (edge-preserving blur)
+// https://people.csail.mit.edu/sparis/bf_course/course_notes.pdf
+//
+// Note: Does not use the Gaussian kernel part of a typical bilateral blur
+// From the paper: "use the information gathered on a neighborhood of 4 x 4 using a bilateral filter for
+// reconstruction, using _uniform_ convolution weights"
+//
+// Note: The paper does a 4x4 (not quite centered) filter, offset by +/- 1 pixel every other frame
+// XeGTAO does a 3x3 filter, on two pixels at a time per compute thread, applied twice
+// We do a 3x3 filter, on 1 pixel per compute thread, applied once
+//
+// Ported from Bevy Engine, crates/bevy_pbr/src/ssao/spatial_denoise.wgsl (v0.13.2), licensed
+// MIT OR Apache-2.0 (see res/licenses/), itself derived from Intel XeGTAO (MIT).
+//
+// PORT: the textureGather calls are rewritten as explicit per-neighbor textureLoads (r32float
+// and r32uint are unfilterable); Bevy view uniforms -> the mod's uniform block; r16float -> r32float.
+
+struct Uniforms {
+    projection: mat4x4f,
+    inverse_projection: mat4x4f,
+    size: vec2f,
+    inv_size: vec2f,
+    depth_scale: vec2f,
+    effect_radius: f32,
+    intensity: f32,
+    slice_count: f32,
+    samples_per_slice_side: f32,
+    debug_view: u32,
+    _pad: f32,
+}
+
+@group(0) @binding(0) var ambient_occlusion_noisy: texture_2d<f32>;
+@group(0) @binding(1) var depth_differences: texture_2d<u32>;
+@group(0) @binding(2) var ambient_occlusion: texture_storage_2d<r32float, write>;
+@group(0) @binding(3) var<uniform> uniforms: Uniforms;
+
+fn clamp_coordinates(pixel_coordinates: vec2<i32>) -> vec2<i32> {
+    return clamp(pixel_coordinates, vec2<i32>(0i), vec2<i32>(uniforms.size) - 1i);
+}
+
+// Each pixel's packed edge info is (left, right, top, bottom) weights, packed by the GTAO pass.
+fn load_edges(pixel_coordinates: vec2<i32>) -> vec4<f32> {
+    return unpack4x8unorm(textureLoad(depth_differences, clamp_coordinates(pixel_coordinates), 0i).r);
+}
+
+fn load_visibility(pixel_coordinates: vec2<i32>) -> f32 {
+    return textureLoad(ambient_occlusion_noisy, clamp_coordinates(pixel_coordinates), 0i).r;
+}
+
+@compute
+@workgroup_size(8, 8, 1)
+fn spatial_denoise(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let pixel_coordinates = vec2<i32>(global_id.xy);
+
+    let left_edges = load_edges(pixel_coordinates + vec2<i32>(-1i, 0i));
+    let right_edges = load_edges(pixel_coordinates + vec2<i32>(1i, 0i));
+    let top_edges = load_edges(pixel_coordinates + vec2<i32>(0i, -1i));
+    let bottom_edges = load_edges(pixel_coordinates + vec2<i32>(0i, 1i));
+    var center_edges = load_edges(pixel_coordinates);
+    // Cross-check each edge against the neighbor's opposing edge weight.
+    center_edges *= vec4<f32>(left_edges.y, right_edges.x, top_edges.w, bottom_edges.z);
+
+    let center_weight = 1.2;
+    let left_weight = center_edges.x;
+    let right_weight = center_edges.y;
+    let top_weight = center_edges.z;
+    let bottom_weight = center_edges.w;
+    let top_left_weight = 0.425 * (top_weight * top_edges.x + left_weight * left_edges.z);
+    let top_right_weight = 0.425 * (top_weight * top_edges.y + right_weight * right_edges.z);
+    let bottom_left_weight = 0.425 * (bottom_weight * bottom_edges.x + left_weight * left_edges.w);
+    let bottom_right_weight = 0.425 * (bottom_weight * bottom_edges.y + right_weight * right_edges.w);
+
+    let center_visibility = load_visibility(pixel_coordinates);
+    let left_visibility = load_visibility(pixel_coordinates + vec2<i32>(-1i, 0i));
+    let right_visibility = load_visibility(pixel_coordinates + vec2<i32>(1i, 0i));
+    let top_visibility = load_visibility(pixel_coordinates + vec2<i32>(0i, -1i));
+    let bottom_visibility = load_visibility(pixel_coordinates + vec2<i32>(0i, 1i));
+    let top_left_visibility = load_visibility(pixel_coordinates + vec2<i32>(-1i, -1i));
+    let top_right_visibility = load_visibility(pixel_coordinates + vec2<i32>(1i, -1i));
+    let bottom_left_visibility = load_visibility(pixel_coordinates + vec2<i32>(-1i, 1i));
+    let bottom_right_visibility = load_visibility(pixel_coordinates + vec2<i32>(1i, 1i));
+
+    // PORT: Bevy sums the center sample unweighted while still counting center_weight in the
+    // denominator; XeGTAO's original weights the value too, which is what we do here.
+    var sum = center_visibility * center_weight;
+    sum += left_visibility * left_weight;
+    sum += right_visibility * right_weight;
+    sum += top_visibility * top_weight;
+    sum += bottom_visibility * bottom_weight;
+    sum += top_left_visibility * top_left_weight;
+    sum += top_right_visibility * top_right_weight;
+    sum += bottom_left_visibility * bottom_left_weight;
+    sum += bottom_right_visibility * bottom_right_weight;
+
+    var sum_weight = center_weight;
+    sum_weight += left_weight;
+    sum_weight += right_weight;
+    sum_weight += top_weight;
+    sum_weight += bottom_weight;
+    sum_weight += top_left_weight;
+    sum_weight += top_right_weight;
+    sum_weight += bottom_left_weight;
+    sum_weight += bottom_right_weight;
+
+    let denoised_visibility = sum / sum_weight;
+
+    textureStore(ambient_occlusion, pixel_coordinates, vec4<f32>(denoised_visibility, 0.0, 0.0, 0.0));
+}

mods/ao_mod/res/gtao.wgsl

@@ -0,0 +1,247 @@
+// Ground Truth-based Ambient Occlusion (GTAO)
+// Paper: https://www.activision.com/cdn/research/Practical_Real_Time_Strategies_for_Accurate_Indirect_Occlusion_NEW%20VERSION_COLOR.pdf
+// Presentation: https://blog.selfshadow.com/publications/s2016-shading-course/activision/s2016_pbs_activision_occlusion.pdf
+//
+// Ported from Bevy Engine, crates/bevy_pbr/src/ssao/gtao.wgsl (v0.13.2), licensed
+// MIT OR Apache-2.0 (see res/licenses/), itself heavily based on XeGTAO v1.30 from Intel (MIT):
+// https://github.com/GameTechDev/XeGTAO/blob/0d177ce06bfa642f64d8af4de1197ad1bcb862d4/Source/Rendering/Shaders/XeGTAO.hlsli
+//
+// PORT:
+// - Bevy view/globals bindings -> the mod's own uniform block (matrices from Dusklight's
+//   CameraService, WebGPU clip convention, reversed-Z - the same convention Bevy uses).
+// - Prepass normals -> normals reconstructed from depth (atyuwen's accurate 5-tap method,
+//   https://atyuwen.github.io/posts/normal-reconstruction/).
+// - Sampler-based reads -> textureLoad (r32float is unfilterable without optional features);
+//   the mip level for the XeGTAO bandwidth optimization is selected explicitly per load.
+// - effect_radius and slice/sample counts come from uniforms instead of constants/shader defs
+//   (game world units are ~100x larger than Bevy's meters, and quality is a live setting).
+// - No TEMPORAL_JITTER: the noise index is pinned (no TAA; the spatial denoiser is the only
+//   filter, a configuration XeGTAO supports).
+// - Storage format r16float -> r32float (core WebGPU storage format).
+
+struct Uniforms {
+    projection: mat4x4f,
+    inverse_projection: mat4x4f,
+    size: vec2f,
+    inv_size: vec2f,
+    depth_scale: vec2f,
+    effect_radius: f32,
+    intensity: f32,
+    slice_count: f32,
+    samples_per_slice_side: f32,
+    debug_view: u32,
+    _pad: f32,
+}
+
+@group(0) @binding(0) var preprocessed_depth: texture_2d<f32>;
+@group(0) @binding(1) var hilbert_index_lut: texture_2d<u32>;
+@group(0) @binding(2) var ambient_occlusion: texture_storage_2d<r32float, write>;
+@group(0) @binding(3) var depth_differences: texture_storage_2d<r32uint, write>;
+@group(0) @binding(4) var<uniform> uniforms: Uniforms;
+
+const PI: f32 = 3.141592653589793;
+const HALF_PI: f32 = 1.5707963267948966;
+
+fn fast_sqrt(x: f32) -> f32 {
+    return bitcast<f32>(0x1fbd1df5 + (bitcast<i32>(x) >> 1u));
+}
+
+fn fast_acos(in_x: f32) -> f32 {
+    let x = abs(in_x);
+    var res = -0.156583 * x + HALF_PI;
+    res *= fast_sqrt(1.0 - x);
+    return select(PI - res, res, in_x >= 0.0);
+}
+
+fn load_noise(pixel_coordinates: vec2<i32>) -> vec2<f32> {
+    let index = textureLoad(hilbert_index_lut, pixel_coordinates % 64, 0).r;
+    // R2 sequence - http://extremelearning.com.au/unreasonable-effectiveness-of-quasirandom-sequences
+    return fract(0.5 + f32(index) * vec2<f32>(0.75487766624669276005, 0.5698402909980532659114));
+}
+
+fn load_depth(pixel_coordinates: vec2<i32>, mip_level: i32) -> f32 {
+    let mip_size = max(vec2<i32>(uniforms.size) >> vec2<u32>(u32(mip_level)), vec2<i32>(1i));
+    let coordinates = clamp(pixel_coordinates, vec2<i32>(0i), mip_size - 1i);
+    return textureLoad(preprocessed_depth, coordinates, mip_level).r;
+}
+
+// Calculate differences in depth between neighbor pixels (later used by the spatial denoiser pass to preserve object edges)
+fn calculate_neighboring_depth_differences(pixel_coordinates: vec2<i32>) -> f32 {
+    // Sample the pixel's depth and 4 depths around it
+    // PORT: explicit loads instead of two textureGathers.
+    let depth_center = load_depth(pixel_coordinates, 0i);
+    let depth_left = load_depth(pixel_coordinates + vec2<i32>(-1i, 0i), 0i);
+    let depth_top = load_depth(pixel_coordinates + vec2<i32>(0i, -1i), 0i);
+    let depth_bottom = load_depth(pixel_coordinates + vec2<i32>(0i, 1i), 0i);
+    let depth_right = load_depth(pixel_coordinates + vec2<i32>(1i, 0i), 0i);
+
+    // Calculate the depth differences (large differences represent object edges)
+    var edge_info = vec4<f32>(depth_left, depth_right, depth_top, depth_bottom) - depth_center;
+    let slope_left_right = (edge_info.y - edge_info.x) * 0.5;
+    let slope_top_bottom = (edge_info.w - edge_info.z) * 0.5;
+    let edge_info_slope_adjusted = edge_info + vec4<f32>(slope_left_right, -slope_left_right, slope_top_bottom, -slope_top_bottom);
+    edge_info = min(abs(edge_info), abs(edge_info_slope_adjusted));
+    let bias = 0.25; // Using the bias and then saturating nudges the values a bit
+    let scale = depth_center * 0.011; // Weight the edges by their distance from the camera
+    edge_info = saturate((1.0 + bias) - edge_info / scale); // Apply the bias and scale, and invert edge_info so that small values become large, and vice versa
+
+    // Pack the edge info into the texture
+    let edge_info_packed = vec4<u32>(pack4x8unorm(edge_info), 0u, 0u, 0u);
+    textureStore(depth_differences, pixel_coordinates, edge_info_packed);
+
+    return depth_center;
+}
+
+fn reconstruct_view_space_position(depth: f32, uv: vec2<f32>) -> vec3<f32> {
+    let clip_xy = vec2<f32>(uv.x * 2.0 - 1.0, 1.0 - 2.0 * uv.y);
+    let t = uniforms.inverse_projection * vec4<f32>(clip_xy, depth, 1.0);
+    let view_xyz = t.xyz / t.w;
+    return view_xyz;
+}
+
+fn view_position_at(pixel_coordinates: vec2<i32>) -> vec3<f32> {
+    let depth = load_depth(pixel_coordinates, 0i);
+    let uv = (vec2<f32>(pixel_coordinates) + 0.5) * uniforms.inv_size;
+    return reconstruct_view_space_position(depth, uv);
+}
+
+// PORT: replaces Bevy's load_normal_view_space (which reads a prepass normal texture we do
+// not have). Accurate view-space normal reconstruction from depth, atyuwen's 5-tap method:
+// for each axis, extrapolate the center depth from the two taps on each side and derive the
+// tangent from whichever side predicts it better. This keeps normals stable across depth
+// discontinuities where naive derivatives smear.
+fn reconstruct_normal(pixel_coordinates: vec2<i32>, pixel_position: vec3<f32>, depth_center: f32) -> vec3<f32> {
+    let depth_left1 = load_depth(pixel_coordinates + vec2<i32>(-1i, 0i), 0i);
+    let depth_left2 = load_depth(pixel_coordinates + vec2<i32>(-2i, 0i), 0i);
+    let depth_right1 = load_depth(pixel_coordinates + vec2<i32>(1i, 0i), 0i);
+    let depth_right2 = load_depth(pixel_coordinates + vec2<i32>(2i, 0i), 0i);
+    let depth_top1 = load_depth(pixel_coordinates + vec2<i32>(0i, -1i), 0i);
+    let depth_top2 = load_depth(pixel_coordinates + vec2<i32>(0i, -2i), 0i);
+    let depth_bottom1 = load_depth(pixel_coordinates + vec2<i32>(0i, 1i), 0i);
+    let depth_bottom2 = load_depth(pixel_coordinates + vec2<i32>(0i, 2i), 0i);
+
+    let use_left = abs(2.0 * depth_left1 - depth_left2 - depth_center) <
+        abs(2.0 * depth_right1 - depth_right2 - depth_center);
+    let use_top = abs(2.0 * depth_top1 - depth_top2 - depth_center) <
+        abs(2.0 * depth_bottom1 - depth_bottom2 - depth_center);
+
+    var ddx: vec3<f32>;
+    if use_left {
+        ddx = pixel_position - view_position_at(pixel_coordinates + vec2<i32>(-1i, 0i));
+    } else {
+        ddx = view_position_at(pixel_coordinates + vec2<i32>(1i, 0i)) - pixel_position;
+    }
+    var ddy: vec3<f32>;
+    if use_top {
+        ddy = pixel_position - view_position_at(pixel_coordinates + vec2<i32>(0i, -1i));
+    } else {
+        ddy = view_position_at(pixel_coordinates + vec2<i32>(0i, 1i)) - pixel_position;
+    }
+
+    var normal = normalize(cross(ddy, ddx));
+    // Guard the orientation: the normal must face the camera.
+    if dot(normal, pixel_position) > 0.0 {
+        normal = -normal;
+    }
+    return normal;
+}
+
+fn load_and_reconstruct_view_space_position(uv: vec2<f32>, sample_mip_level: f32) -> vec3<f32> {
+    // PORT: point-sample the selected mip explicitly instead of textureSampleLevel.
+    let mip_level = i32(sample_mip_level + 0.5);
+    let mip_size = max(vec2<i32>(uniforms.size) >> vec2<u32>(u32(mip_level)), vec2<i32>(1i));
+    let depth = load_depth(vec2<i32>(uv * vec2<f32>(mip_size)), mip_level);
+    return reconstruct_view_space_position(depth, uv);
+}
+
+@compute
+@workgroup_size(8, 8, 1)
+fn gtao(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let slice_count = uniforms.slice_count;
+    let samples_per_slice_side = uniforms.samples_per_slice_side;
+    let effect_radius = uniforms.effect_radius;
+    let falloff_range = 0.615 * effect_radius;
+    let falloff_from = effect_radius * (1.0 - 0.615);
+    let falloff_mul = -1.0 / falloff_range;
+    let falloff_add = falloff_from / falloff_range + 1.0;
+
+    let pixel_coordinates = vec2<i32>(global_id.xy);
+    let uv = (vec2<f32>(pixel_coordinates) + 0.5) * uniforms.inv_size;
+
+    var pixel_depth = calculate_neighboring_depth_differences(pixel_coordinates);
+    let raw_depth = pixel_depth;
+    pixel_depth += 0.00001; // Avoid depth precision issues
+
+    let pixel_position = reconstruct_view_space_position(pixel_depth, uv);
+    // PORT: the reconstruction differences the center position against neighbor positions
+    // built from unbiased depths, so its center must use the raw depth too: at this game's
+    // depth scale (far plane 200000 -> depth ~5e-3) Bevy's +0.00001 bias is comparable to a
+    // one-pixel depth step, and a biased center corrupts both tangents.
+    let pixel_normal = reconstruct_normal(
+        pixel_coordinates, reconstruct_view_space_position(raw_depth, uv), raw_depth);
+    let view_vec = normalize(-pixel_position);
+
+    let noise = load_noise(pixel_coordinates);
+    let sample_scale = (-0.5 * effect_radius * uniforms.projection[0][0]) / pixel_position.z;
+
+    var visibility = 0.0;
+    for (var slice_t = 0.0; slice_t < slice_count; slice_t += 1.0) {
+        let slice = slice_t + noise.x;
+        let phi = (PI / slice_count) * slice;
+        let omega = vec2<f32>(cos(phi), sin(phi));
+
+        let direction = vec3<f32>(omega.xy, 0.0);
+        let orthographic_direction = direction - (dot(direction, view_vec) * view_vec);
+        let axis = cross(direction, view_vec);
+        let projected_normal = pixel_normal - axis * dot(pixel_normal, axis);
+        let projected_normal_length = length(projected_normal);
+
+        let sign_norm = sign(dot(orthographic_direction, projected_normal));
+        let cos_norm = saturate(dot(projected_normal, view_vec) / projected_normal_length);
+        let n = sign_norm * fast_acos(cos_norm);
+
+        let min_cos_horizon_1 = cos(n + HALF_PI);
+        let min_cos_horizon_2 = cos(n - HALF_PI);
+        var cos_horizon_1 = min_cos_horizon_1;
+        var cos_horizon_2 = min_cos_horizon_2;
+        let sample_mul = vec2<f32>(omega.x, -omega.y) * sample_scale;
+        for (var sample_t = 0.0; sample_t < samples_per_slice_side; sample_t += 1.0) {
+            var sample_noise = (slice_t + sample_t * samples_per_slice_side) * 0.6180339887498948482;
+            sample_noise = fract(noise.y + sample_noise);
+
+            var s = (sample_t + sample_noise) / samples_per_slice_side;
+            s *= s; // https://github.com/GameTechDev/XeGTAO#sample-distribution
+            let sample = s * sample_mul;
+
+            // * uniforms.size gets us from [0, 1] to [0, viewport_size], which is needed for this to get the correct mip levels
+            let sample_mip_level = clamp(log2(length(sample * uniforms.size)) - 3.3, 0.0, 4.0); // https://github.com/GameTechDev/XeGTAO#memory-bandwidth-bottleneck
+            let sample_position_1 = load_and_reconstruct_view_space_position(uv + sample, sample_mip_level);
+            let sample_position_2 = load_and_reconstruct_view_space_position(uv - sample, sample_mip_level);
+
+            let sample_difference_1 = sample_position_1 - pixel_position;
+            let sample_difference_2 = sample_position_2 - pixel_position;
+            let sample_distance_1 = length(sample_difference_1);
+            let sample_distance_2 = length(sample_difference_2);
+            var sample_cos_horizon_1 = dot(sample_difference_1 / sample_distance_1, view_vec);
+            var sample_cos_horizon_2 = dot(sample_difference_2 / sample_distance_2, view_vec);
+
+            let weight_1 = saturate(sample_distance_1 * falloff_mul + falloff_add);
+            let weight_2 = saturate(sample_distance_2 * falloff_mul + falloff_add);
+            sample_cos_horizon_1 = mix(min_cos_horizon_1, sample_cos_horizon_1, weight_1);
+            sample_cos_horizon_2 = mix(min_cos_horizon_2, sample_cos_horizon_2, weight_2);
+
+            cos_horizon_1 = max(cos_horizon_1, sample_cos_horizon_1);
+            cos_horizon_2 = max(cos_horizon_2, sample_cos_horizon_2);
+        }
+
+        let horizon_1 = fast_acos(cos_horizon_1);
+        let horizon_2 = -fast_acos(cos_horizon_2);
+        let v1 = (cos_norm + 2.0 * horizon_1 * sin(n) - cos(2.0 * horizon_1 - n)) / 4.0;
+        let v2 = (cos_norm + 2.0 * horizon_2 * sin(n) - cos(2.0 * horizon_2 - n)) / 4.0;
+        visibility += projected_normal_length * (v1 + v2);
+    }
+    visibility /= slice_count;
+    visibility = clamp(visibility, 0.03, 1.0);
+
+    textureStore(ambient_occlusion, pixel_coordinates, vec4<f32>(visibility, 0.0, 0.0, 0.0));
+}

mods/ao_mod/res/licenses/BEVY-APACHE-2.0.txt

@@ -0,0 +1,176 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS

mods/ao_mod/res/licenses/BEVY-MIT.txt

@@ -0,0 +1,19 @@
+MIT License
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

mods/ao_mod/res/licenses/XEGTAO-MIT.txt

@@ -0,0 +1,21 @@
+MIT License
+
+Copyright (C) 2016-2021, Intel Corporation 
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

mods/ao_mod/res/preprocess_depth.wgsl

@@ -0,0 +1,138 @@
+// Inputs a depth texture and outputs a MIP-chain of depths.
+//
+// Because SSAO's performance is bound by texture reads, this increases
+// performance over using the full resolution depth for every sample.
+//
+// Reference: https://research.nvidia.com/sites/default/files/pubs/2012-06_Scalable-Ambient-Obscurance/McGuire12SAO.pdf, section 2.2
+//
+// Ported from Bevy Engine, crates/bevy_pbr/src/ssao/preprocess_depth.wgsl (v0.13.2),
+// licensed MIT OR Apache-2.0 (see res/licenses/), itself derived from Intel XeGTAO (MIT).
+//
+// PORT: sampler-based gathers replaced with textureLoad (r32float is not filterable without
+// optional device features), Bevy view uniforms replaced with the mod's own uniform block,
+// storage format r16float -> r32float (core WebGPU storage format). MIP 4 moved into its own
+// entry point (core WebGPU limit is 4 storage textures per stage).
+
+struct Uniforms {
+    projection: mat4x4f,
+    inverse_projection: mat4x4f,
+    size: vec2f,        // AO chain size in pixels (MIP 0 of the preprocessed depth)
+    inv_size: vec2f,
+    depth_scale: vec2f, // input depth snapshot pixels per chain pixel (1 or 2)
+    effect_radius: f32, // view-space units
+    intensity: f32,
+    slice_count: f32,
+    samples_per_slice_side: f32,
+    debug_view: u32,
+    _pad: f32,
+}
+
+@group(0) @binding(0) var input_depth: texture_2d<f32>;
+@group(0) @binding(1) var preprocessed_depth_mip0: texture_storage_2d<r32float, write>;
+@group(0) @binding(2) var preprocessed_depth_mip1: texture_storage_2d<r32float, write>;
+@group(0) @binding(3) var preprocessed_depth_mip2: texture_storage_2d<r32float, write>;
+@group(0) @binding(4) var preprocessed_depth_mip3: texture_storage_2d<r32float, write>;
+@group(0) @binding(5) var<uniform> uniforms: Uniforms;
+// downsample_mip4 entry point only (disjoint subresources of the same texture).
+@group(0) @binding(6) var preprocessed_depth_mip3_in: texture_2d<f32>;
+@group(0) @binding(7) var preprocessed_depth_mip4: texture_storage_2d<r32float, write>;
+
+// PORT: replaces the textureGather of the input depth with explicit loads (also handles the
+// half-resolution case, where one chain texel covers depth_scale snapshot texels).
+fn load_input_depth(pixel_coordinates: vec2<i32>) -> f32 {
+    let input_size = vec2<i32>(uniforms.size * uniforms.depth_scale);
+    let coordinates = clamp(vec2<i32>(vec2<f32>(pixel_coordinates) * uniforms.depth_scale),
+        vec2<i32>(0i), input_size - 1i);
+    return textureLoad(input_depth, coordinates, 0i).r;
+}
+
+// Using 4 depths from the previous MIP, compute a weighted average for the depth of the current MIP
+fn weighted_average(depth0: f32, depth1: f32, depth2: f32, depth3: f32) -> f32 {
+    let depth_range_scale_factor = 0.75;
+    let effect_radius = depth_range_scale_factor * 0.5 * 1.457;
+    let falloff_range = 0.615 * effect_radius;
+    let falloff_from = effect_radius * (1.0 - 0.615);
+    let falloff_mul = -1.0 / falloff_range;
+    let falloff_add = falloff_from / falloff_range + 1.0;
+
+    let min_depth = min(min(depth0, depth1), min(depth2, depth3));
+    let weight0 = saturate((depth0 - min_depth) * falloff_mul + falloff_add);
+    let weight1 = saturate((depth1 - min_depth) * falloff_mul + falloff_add);
+    let weight2 = saturate((depth2 - min_depth) * falloff_mul + falloff_add);
+    let weight3 = saturate((depth3 - min_depth) * falloff_mul + falloff_add);
+    let weight_total = weight0 + weight1 + weight2 + weight3;
+
+    return ((weight0 * depth0) + (weight1 * depth1) + (weight2 * depth2) + (weight3 * depth3)) / weight_total;
+}
+
+// Used to share the depths from the previous MIP level between all invocations in a workgroup
+var<workgroup> previous_mip_depth: array<array<f32, 8>, 8>;
+
+@compute
+@workgroup_size(8, 8, 1)
+fn preprocess_depth(@builtin(global_invocation_id) global_id: vec3<u32>, @builtin(local_invocation_id) local_id: vec3<u32>) {
+    let base_coordinates = vec2<i32>(global_id.xy);
+
+    // MIP 0 - Copy 4 texels from the input depth (per invocation, 8x8 invocations per workgroup)
+    let pixel_coordinates0 = base_coordinates * 2i;
+    let pixel_coordinates1 = pixel_coordinates0 + vec2<i32>(1i, 0i);
+    let pixel_coordinates2 = pixel_coordinates0 + vec2<i32>(0i, 1i);
+    let pixel_coordinates3 = pixel_coordinates0 + vec2<i32>(1i, 1i);
+    let depth0 = load_input_depth(pixel_coordinates0);
+    let depth1 = load_input_depth(pixel_coordinates1);
+    let depth2 = load_input_depth(pixel_coordinates2);
+    let depth3 = load_input_depth(pixel_coordinates3);
+    textureStore(preprocessed_depth_mip0, pixel_coordinates0, vec4<f32>(depth0, 0.0, 0.0, 0.0));
+    textureStore(preprocessed_depth_mip0, pixel_coordinates1, vec4<f32>(depth1, 0.0, 0.0, 0.0));
+    textureStore(preprocessed_depth_mip0, pixel_coordinates2, vec4<f32>(depth2, 0.0, 0.0, 0.0));
+    textureStore(preprocessed_depth_mip0, pixel_coordinates3, vec4<f32>(depth3, 0.0, 0.0, 0.0));
+
+    // MIP 1 - Weighted average of MIP 0's depth values (per invocation, 8x8 invocations per workgroup)
+    let depth_mip1 = weighted_average(depth0, depth1, depth2, depth3);
+    textureStore(preprocessed_depth_mip1, base_coordinates, vec4<f32>(depth_mip1, 0.0, 0.0, 0.0));
+    previous_mip_depth[local_id.x][local_id.y] = depth_mip1;
+
+    workgroupBarrier();
+
+    // MIP 2 - Weighted average of MIP 1's depth values (per invocation, 4x4 invocations per workgroup)
+    if all(local_id.xy % vec2<u32>(2u) == vec2<u32>(0u)) {
+        let mip2_depth0 = previous_mip_depth[local_id.x + 0u][local_id.y + 0u];
+        let mip2_depth1 = previous_mip_depth[local_id.x + 1u][local_id.y + 0u];
+        let mip2_depth2 = previous_mip_depth[local_id.x + 0u][local_id.y + 1u];
+        let mip2_depth3 = previous_mip_depth[local_id.x + 1u][local_id.y + 1u];
+        let depth_mip2 = weighted_average(mip2_depth0, mip2_depth1, mip2_depth2, mip2_depth3);
+        textureStore(preprocessed_depth_mip2, base_coordinates / 2i, vec4<f32>(depth_mip2, 0.0, 0.0, 0.0));
+        previous_mip_depth[local_id.x][local_id.y] = depth_mip2;
+    }
+
+    workgroupBarrier();
+
+    // MIP 3 - Weighted average of MIP 2's depth values (per invocation, 2x2 invocations per workgroup)
+    if all(local_id.xy % vec2<u32>(4u) == vec2<u32>(0u)) {
+        let mip3_depth0 = previous_mip_depth[local_id.x + 0u][local_id.y + 0u];
+        let mip3_depth1 = previous_mip_depth[local_id.x + 2u][local_id.y + 0u];
+        let mip3_depth2 = previous_mip_depth[local_id.x + 0u][local_id.y + 2u];
+        let mip3_depth3 = previous_mip_depth[local_id.x + 2u][local_id.y + 2u];
+        let depth_mip3 = weighted_average(mip3_depth0, mip3_depth1, mip3_depth2, mip3_depth3);
+        textureStore(preprocessed_depth_mip3, base_coordinates / 4i, vec4<f32>(depth_mip3, 0.0, 0.0, 0.0));
+        previous_mip_depth[local_id.x][local_id.y] = depth_mip3;
+    }
+}
+
+// MIP 4: weighted average of MIP 3's depth values, as a second (tiny) dispatch.
+@compute
+@workgroup_size(8, 8, 1)
+fn downsample_mip4(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    let base_coordinates = vec2<i32>(global_id.xy);
+    let mip3_size = max(vec2<i32>(textureDimensions(preprocessed_depth_mip3_in)), vec2<i32>(1i));
+    let coordinates0 = clamp(base_coordinates * 2i, vec2<i32>(0i), mip3_size - 1i);
+    let coordinates1 = clamp(base_coordinates * 2i + vec2<i32>(1i, 0i), vec2<i32>(0i), mip3_size - 1i);
+    let coordinates2 = clamp(base_coordinates * 2i + vec2<i32>(0i, 1i), vec2<i32>(0i), mip3_size - 1i);
+    let coordinates3 = clamp(base_coordinates * 2i + vec2<i32>(1i, 1i), vec2<i32>(0i), mip3_size - 1i);
+    let depth0 = textureLoad(preprocessed_depth_mip3_in, coordinates0, 0i).r;
+    let depth1 = textureLoad(preprocessed_depth_mip3_in, coordinates1, 0i).r;
+    let depth2 = textureLoad(preprocessed_depth_mip3_in, coordinates2, 0i).r;
+    let depth3 = textureLoad(preprocessed_depth_mip3_in, coordinates3, 0i).r;
+    let depth_mip4 = weighted_average(depth0, depth1, depth2, depth3);
+    textureStore(preprocessed_depth_mip4, base_coordinates, vec4<f32>(depth_mip4, 0.0, 0.0, 0.0));
+}

mods/template_mod/mod.json

@@ -0,0 +1,7 @@
+{
+  "id": "com.example.mod",
+  "name": "Template Mod",
+  "version": "1.0.0",
+  "author": "You",
+  "description": "An example Dusklight mod"
+}

sdk/CMakeLists.txt

@@ -1,7 +1,7 @@
 # Dusklight Mod SDK entry point
 #
-# Provides game/service headers, compile definitions and version.h without
-# configuring the full game tree.
+# Provides game/service headers, compile definitions, version.h and WebGPU
+# headers without configuring the full game tree.
 #
 # Usage (from a mod project):
 #   add_subdirectory(<dusk>/sdk dusk-sdk EXCLUDE_FROM_ALL)
@@ -30,6 +30,11 @@ include("${CMAKE_CURRENT_SOURCE_DIR}/../cmake/DetectVersion.cmake")
 detect_version()
 configure_version_header()
 
+# Provides dawn::webgpu_dawn and dawn::dawncpp_headers for public gfx service headers.
+include("${CMAKE_CURRENT_SOURCE_DIR}/../extern/aurora/cmake/AuroraDependencyVersions.cmake")
+set(AURORA_DAWN_PROVIDER "package" CACHE STRING "How to provide Dawn for the mod SDK")
+include("${CMAKE_CURRENT_SOURCE_DIR}/../extern/aurora/cmake/AuroraDawnProvider.cmake")
+
 # Game ABI headers & compile definitions
 include("${CMAKE_CURRENT_SOURCE_DIR}/../cmake/GameABIConfig.cmake")
 

tools/mod_template/mod.json

@@ -1,7 +0,0 @@
-{
-    "id": "com.example.mod",
-    "name": "Template Mod",
-    "version": "1.0.0",
-    "author": "You",
-    "description": "An example Dusklight mod"
-}