From 75a249f3e13858f6c70133763ff7c075c46a34bb Mon Sep 17 00:00:00 2001 From: Luke Street Date: Fri, 10 Jul 2026 00:40:18 -0600 Subject: [PATCH] GTAO demo mod --- CMakeLists.txt | 3 +- cmake/GameABIConfig.cmake | 5 + docs/modding.md | 2 +- extern/aurora | 2 +- mods/ao_mod/CMakeLists.txt | 20 + mods/ao_mod/mod.json | 7 + mods/ao_mod/res/composite.wgsl | 161 +++ mods/ao_mod/res/denoise.wgsl | 108 ++ mods/ao_mod/res/gtao.wgsl | 247 +++++ mods/ao_mod/res/licenses/BEVY-APACHE-2.0.txt | 176 ++++ mods/ao_mod/res/licenses/BEVY-MIT.txt | 19 + mods/ao_mod/res/licenses/XEGTAO-MIT.txt | 21 + mods/ao_mod/res/preprocess_depth.wgsl | 138 +++ mods/ao_mod/src/mod.cpp | 931 ++++++++++++++++++ .../template_mod}/CMakeLists.txt | 0 mods/template_mod/mod.json | 7 + .../template_mod}/res/.gitkeep | 0 .../template_mod}/src/mod.cpp | 0 sdk/CMakeLists.txt | 9 +- tools/mod_template/mod.json | 7 - 20 files changed, 1851 insertions(+), 12 deletions(-) create mode 100644 mods/ao_mod/CMakeLists.txt create mode 100644 mods/ao_mod/mod.json create mode 100644 mods/ao_mod/res/composite.wgsl create mode 100644 mods/ao_mod/res/denoise.wgsl create mode 100644 mods/ao_mod/res/gtao.wgsl create mode 100644 mods/ao_mod/res/licenses/BEVY-APACHE-2.0.txt create mode 100644 mods/ao_mod/res/licenses/BEVY-MIT.txt create mode 100644 mods/ao_mod/res/licenses/XEGTAO-MIT.txt create mode 100644 mods/ao_mod/res/preprocess_depth.wgsl create mode 100644 mods/ao_mod/src/mod.cpp rename {tools/mod_template => mods/template_mod}/CMakeLists.txt (100%) create mode 100644 mods/template_mod/mod.json rename {tools/mod_template => mods/template_mod}/res/.gitkeep (100%) rename {tools/mod_template => mods/template_mod}/src/mod.cpp (100%) delete mode 100644 tools/mod_template/mod.json diff --git a/CMakeLists.txt b/CMakeLists.txt index e62704fe1c..abf30f6863 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/cmake/GameABIConfig.cmake b/cmake/GameABIConfig.cmake index 030f1181a0..d48be445a2 100644 --- a/cmake/GameABIConfig.cmake +++ b/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 () diff --git a/docs/modding.md b/docs/modding.md index b2e7ba2e26..6f2d1a8580 100644 --- a/docs/modding.md +++ b/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/ diff --git a/extern/aurora b/extern/aurora index 0a5a5d90ef..1dde08fa0d 160000 --- a/extern/aurora +++ b/extern/aurora @@ -1 +1 @@ -Subproject commit 0a5a5d90efd2dbc7b402ab61f1021922dbf7496f +Subproject commit 1dde08fa0d0030133788a6250a81c8b9c44f246f diff --git a/mods/ao_mod/CMakeLists.txt b/mods/ao_mod/CMakeLists.txt new file mode 100644 index 0000000000..13590923f0 --- /dev/null +++ b/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 +) diff --git a/mods/ao_mod/mod.json b/mods/ao_mod/mod.json new file mode 100644 index 0000000000..9019cf335d --- /dev/null +++ b/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." +} diff --git a/mods/ao_mod/res/composite.wgsl b/mods/ao_mod/res/composite.wgsl new file mode 100644 index 0000000000..fa3a0101b8 --- /dev/null +++ b/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; +@group(0) @binding(1) var preprocessed_depth: texture_2d; +@group(0) @binding(2) var scene_depth_raw: texture_2d; +@group(0) @binding(3) var 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) -> f32 { + let coordinates = clamp(pixel_coordinates, vec2(0i), vec2(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) -> 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, pixel_position: vec3f, depth_center: f32) -> vec3f { + let depth_left1 = load_depth(pixel_coordinates + vec2(-1i, 0i)); + let depth_left2 = load_depth(pixel_coordinates + vec2(-2i, 0i)); + let depth_right1 = load_depth(pixel_coordinates + vec2(1i, 0i)); + let depth_right2 = load_depth(pixel_coordinates + vec2(2i, 0i)); + let depth_top1 = load_depth(pixel_coordinates + vec2(0i, -1i)); + let depth_top2 = load_depth(pixel_coordinates + vec2(0i, -2i)); + let depth_bottom1 = load_depth(pixel_coordinates + vec2(0i, 1i)); + let depth_bottom2 = load_depth(pixel_coordinates + vec2(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(-1i, 0i)); + } else { + ddx = view_position_at(pixel_coordinates + vec2(1i, 0i)) - pixel_position; + } + var ddy: vec3f; + if use_top { + ddy = pixel_position - view_position_at(pixel_coordinates + vec2(0i, -1i)); + } else { + ddy = view_position_at(pixel_coordinates + vec2(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) -> f32 { + let size = vec2(textureDimensions(scene_depth_raw)); + let coordinates = clamp(pixel_coordinates, vec2(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(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(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(in.uv * size); + let d_center = load_raw_depth(pixel); + let d_left = load_raw_depth(pixel + vec2(-1i, 0i)); + let d_right = load_raw_depth(pixel + vec2(1i, 0i)); + let d_top = load_raw_depth(pixel + vec2(0i, -1i)); + let d_bottom = load_raw_depth(pixel + vec2(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); +} diff --git a/mods/ao_mod/res/denoise.wgsl b/mods/ao_mod/res/denoise.wgsl new file mode 100644 index 0000000000..b947b1bd04 --- /dev/null +++ b/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; +@group(0) @binding(1) var depth_differences: texture_2d; +@group(0) @binding(2) var ambient_occlusion: texture_storage_2d; +@group(0) @binding(3) var uniforms: Uniforms; + +fn clamp_coordinates(pixel_coordinates: vec2) -> vec2 { + return clamp(pixel_coordinates, vec2(0i), vec2(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) -> vec4 { + return unpack4x8unorm(textureLoad(depth_differences, clamp_coordinates(pixel_coordinates), 0i).r); +} + +fn load_visibility(pixel_coordinates: vec2) -> 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) { + let pixel_coordinates = vec2(global_id.xy); + + let left_edges = load_edges(pixel_coordinates + vec2(-1i, 0i)); + let right_edges = load_edges(pixel_coordinates + vec2(1i, 0i)); + let top_edges = load_edges(pixel_coordinates + vec2(0i, -1i)); + let bottom_edges = load_edges(pixel_coordinates + vec2(0i, 1i)); + var center_edges = load_edges(pixel_coordinates); + // Cross-check each edge against the neighbor's opposing edge weight. + center_edges *= vec4(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(-1i, 0i)); + let right_visibility = load_visibility(pixel_coordinates + vec2(1i, 0i)); + let top_visibility = load_visibility(pixel_coordinates + vec2(0i, -1i)); + let bottom_visibility = load_visibility(pixel_coordinates + vec2(0i, 1i)); + let top_left_visibility = load_visibility(pixel_coordinates + vec2(-1i, -1i)); + let top_right_visibility = load_visibility(pixel_coordinates + vec2(1i, -1i)); + let bottom_left_visibility = load_visibility(pixel_coordinates + vec2(-1i, 1i)); + let bottom_right_visibility = load_visibility(pixel_coordinates + vec2(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(denoised_visibility, 0.0, 0.0, 0.0)); +} diff --git a/mods/ao_mod/res/gtao.wgsl b/mods/ao_mod/res/gtao.wgsl new file mode 100644 index 0000000000..ee23774681 --- /dev/null +++ b/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; +@group(0) @binding(1) var hilbert_index_lut: texture_2d; +@group(0) @binding(2) var ambient_occlusion: texture_storage_2d; +@group(0) @binding(3) var depth_differences: texture_storage_2d; +@group(0) @binding(4) var uniforms: Uniforms; + +const PI: f32 = 3.141592653589793; +const HALF_PI: f32 = 1.5707963267948966; + +fn fast_sqrt(x: f32) -> f32 { + return bitcast(0x1fbd1df5 + (bitcast(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) -> vec2 { + 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(0.75487766624669276005, 0.5698402909980532659114)); +} + +fn load_depth(pixel_coordinates: vec2, mip_level: i32) -> f32 { + let mip_size = max(vec2(uniforms.size) >> vec2(u32(mip_level)), vec2(1i)); + let coordinates = clamp(pixel_coordinates, vec2(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) -> 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(-1i, 0i), 0i); + let depth_top = load_depth(pixel_coordinates + vec2(0i, -1i), 0i); + let depth_bottom = load_depth(pixel_coordinates + vec2(0i, 1i), 0i); + let depth_right = load_depth(pixel_coordinates + vec2(1i, 0i), 0i); + + // Calculate the depth differences (large differences represent object edges) + var edge_info = vec4(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(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(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) -> vec3 { + let clip_xy = vec2(uv.x * 2.0 - 1.0, 1.0 - 2.0 * uv.y); + let t = uniforms.inverse_projection * vec4(clip_xy, depth, 1.0); + let view_xyz = t.xyz / t.w; + return view_xyz; +} + +fn view_position_at(pixel_coordinates: vec2) -> vec3 { + let depth = load_depth(pixel_coordinates, 0i); + let uv = (vec2(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, pixel_position: vec3, depth_center: f32) -> vec3 { + let depth_left1 = load_depth(pixel_coordinates + vec2(-1i, 0i), 0i); + let depth_left2 = load_depth(pixel_coordinates + vec2(-2i, 0i), 0i); + let depth_right1 = load_depth(pixel_coordinates + vec2(1i, 0i), 0i); + let depth_right2 = load_depth(pixel_coordinates + vec2(2i, 0i), 0i); + let depth_top1 = load_depth(pixel_coordinates + vec2(0i, -1i), 0i); + let depth_top2 = load_depth(pixel_coordinates + vec2(0i, -2i), 0i); + let depth_bottom1 = load_depth(pixel_coordinates + vec2(0i, 1i), 0i); + let depth_bottom2 = load_depth(pixel_coordinates + vec2(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; + if use_left { + ddx = pixel_position - view_position_at(pixel_coordinates + vec2(-1i, 0i)); + } else { + ddx = view_position_at(pixel_coordinates + vec2(1i, 0i)) - pixel_position; + } + var ddy: vec3; + if use_top { + ddy = pixel_position - view_position_at(pixel_coordinates + vec2(0i, -1i)); + } else { + ddy = view_position_at(pixel_coordinates + vec2(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, sample_mip_level: f32) -> vec3 { + // PORT: point-sample the selected mip explicitly instead of textureSampleLevel. + let mip_level = i32(sample_mip_level + 0.5); + let mip_size = max(vec2(uniforms.size) >> vec2(u32(mip_level)), vec2(1i)); + let depth = load_depth(vec2(uv * vec2(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) { + 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(global_id.xy); + let uv = (vec2(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(cos(phi), sin(phi)); + + let direction = vec3(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(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(visibility, 0.0, 0.0, 0.0)); +} diff --git a/mods/ao_mod/res/licenses/BEVY-APACHE-2.0.txt b/mods/ao_mod/res/licenses/BEVY-APACHE-2.0.txt new file mode 100644 index 0000000000..d9a10c0d8e --- /dev/null +++ b/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 diff --git a/mods/ao_mod/res/licenses/BEVY-MIT.txt b/mods/ao_mod/res/licenses/BEVY-MIT.txt new file mode 100644 index 0000000000..9cf106272a --- /dev/null +++ b/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. diff --git a/mods/ao_mod/res/licenses/XEGTAO-MIT.txt b/mods/ao_mod/res/licenses/XEGTAO-MIT.txt new file mode 100644 index 0000000000..2b1bd1561b --- /dev/null +++ b/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. diff --git a/mods/ao_mod/res/preprocess_depth.wgsl b/mods/ao_mod/res/preprocess_depth.wgsl new file mode 100644 index 0000000000..98fe1d5400 --- /dev/null +++ b/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; +@group(0) @binding(1) var preprocessed_depth_mip0: texture_storage_2d; +@group(0) @binding(2) var preprocessed_depth_mip1: texture_storage_2d; +@group(0) @binding(3) var preprocessed_depth_mip2: texture_storage_2d; +@group(0) @binding(4) var preprocessed_depth_mip3: texture_storage_2d; +@group(0) @binding(5) var uniforms: Uniforms; +// downsample_mip4 entry point only (disjoint subresources of the same texture). +@group(0) @binding(6) var preprocessed_depth_mip3_in: texture_2d; +@group(0) @binding(7) var preprocessed_depth_mip4: texture_storage_2d; + +// 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) -> f32 { + let input_size = vec2(uniforms.size * uniforms.depth_scale); + let coordinates = clamp(vec2(vec2(pixel_coordinates) * uniforms.depth_scale), + vec2(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 previous_mip_depth: array, 8>; + +@compute +@workgroup_size(8, 8, 1) +fn preprocess_depth(@builtin(global_invocation_id) global_id: vec3, @builtin(local_invocation_id) local_id: vec3) { + let base_coordinates = vec2(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(1i, 0i); + let pixel_coordinates2 = pixel_coordinates0 + vec2(0i, 1i); + let pixel_coordinates3 = pixel_coordinates0 + vec2(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(depth0, 0.0, 0.0, 0.0)); + textureStore(preprocessed_depth_mip0, pixel_coordinates1, vec4(depth1, 0.0, 0.0, 0.0)); + textureStore(preprocessed_depth_mip0, pixel_coordinates2, vec4(depth2, 0.0, 0.0, 0.0)); + textureStore(preprocessed_depth_mip0, pixel_coordinates3, vec4(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(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(2u) == vec2(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(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(4u) == vec2(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(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) { + let base_coordinates = vec2(global_id.xy); + let mip3_size = max(vec2(textureDimensions(preprocessed_depth_mip3_in)), vec2(1i)); + let coordinates0 = clamp(base_coordinates * 2i, vec2(0i), mip3_size - 1i); + let coordinates1 = clamp(base_coordinates * 2i + vec2(1i, 0i), vec2(0i), mip3_size - 1i); + let coordinates2 = clamp(base_coordinates * 2i + vec2(0i, 1i), vec2(0i), mip3_size - 1i); + let coordinates3 = clamp(base_coordinates * 2i + vec2(1i, 1i), vec2(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(depth_mip4, 0.0, 0.0, 0.0)); +} diff --git a/mods/ao_mod/src/mod.cpp b/mods/ao_mod/src/mod.cpp new file mode 100644 index 0000000000..a7eb4f4028 --- /dev/null +++ b/mods/ao_mod/src/mod.cpp @@ -0,0 +1,931 @@ +// Ambient occlusion (GTAO) example mod. +// +// Showcases the gfx service's compute tasks and the camera service: after opaque scene draws, +// before translucent/fog overlays, the scene depth is resolved and a three-dispatch compute +// chain (depth MIP prefilter, GTAO, spatial denoise) produces a visibility texture that a +// fullscreen draw multiplies over the world. +// +// The WGSL in res/ is ported from Bevy Engine's SSAO implementation (MIT OR Apache-2.0), +// itself based on Intel XeGTAO (MIT); see res/licenses/ and the `PORT:` notes in the shaders. + +#include "mods/service.hpp" +#include "mods/svc/camera.h" +#include "mods/svc/config.h" +#include "mods/svc/gfx.h" +#include "mods/svc/log.h" +#include "mods/svc/resource.h" +#include "mods/svc/ui.h" + +#include +#include +#include +#include +#include +#include +#include +#include + +DEFINE_MOD(); +IMPORT_SERVICE(LogService, svc_log); +IMPORT_SERVICE(ConfigService, svc_config); +IMPORT_SERVICE(ResourceService, svc_resource); +IMPORT_SERVICE(UiService, svc_ui); +IMPORT_SERVICE(GfxService, svc_gfx); +IMPORT_SERVICE(CameraService, svc_camera); + +namespace { + +ConfigVarHandle g_cvarEnabled = 0; +ConfigVarHandle g_cvarQuality = 0; +ConfigVarHandle g_cvarRadius = 0; +ConfigVarHandle g_cvarIntensity = 0; +ConfigVarHandle g_cvarHalfRes = 0; +ConfigVarHandle g_cvarDebugView = 0; + +GfxComputeTypeHandle g_computeType = 0; +GfxDrawTypeHandle g_drawType = 0; +GfxStageHookHandle g_afterOpaqueHook = 0; +UiWindowHandle g_controlsWindow = 0; + +ResourceBuffer g_preprocessSource = RESOURCE_BUFFER_INIT; +ResourceBuffer g_gtaoSource = RESOURCE_BUFFER_INIT; +ResourceBuffer g_denoiseSource = RESOURCE_BUFFER_INIT; +ResourceBuffer g_compositeSource = RESOURCE_BUFFER_INIT; + +GfxDeviceInfo g_deviceInfo = GFX_DEVICE_INFO_INIT; +WGPUComputePipeline g_preprocessPipeline = nullptr; +WGPUComputePipeline g_mip4Pipeline = nullptr; +WGPUComputePipeline g_gtaoPipeline = nullptr; +WGPUComputePipeline g_denoisePipeline = nullptr; +WGPUBindGroupLayout g_preprocessLayout = nullptr; +WGPUBindGroupLayout g_mip4Layout = nullptr; +WGPUBindGroupLayout g_gtaoLayout = nullptr; +WGPUBindGroupLayout g_denoiseLayout = nullptr; +WGPURenderPipeline g_compositePipeline = nullptr; +WGPURenderPipeline g_compositeDebugPipeline = nullptr; +WGPUBindGroupLayout g_compositeLayout = nullptr; +WGPUBindGroupLayout g_compositeDebugLayout = nullptr; +WGPUTexture g_hilbertLut = nullptr; +WGPUTextureView g_hilbertLutView = nullptr; + +// AO chain targets, recreated when the render size (or halfRes) changes. Old sets are retired +// for a few frames instead of released immediately: payloads embedding their views may still +// be in flight on the render worker. +struct AoTargets { + uint32_t width = 0; + uint32_t height = 0; + WGPUTexture preprocessedDepth = nullptr; + WGPUTextureView preprocessedDepthMips[5] = {}; + WGPUTextureView preprocessedDepthAll = nullptr; + WGPUTexture aoNoisy = nullptr; + WGPUTextureView aoNoisyView = nullptr; + WGPUTexture depthDifferences = nullptr; + WGPUTextureView depthDifferencesView = nullptr; + WGPUTexture aoFinal = nullptr; + WGPUTextureView aoFinalView = nullptr; +}; +AoTargets g_targets; +struct RetiredTargets { + AoTargets targets; + int framesLeft = 0; +}; +std::vector g_retiredTargets; + +bool g_warnedNoDepth = false; +bool g_loggedChain = false; +std::atomic g_chainExecuted{false}; + +// Mirror of the WGSL Uniforms struct (keep in sync with res/*.wgsl). +struct AoUniforms { + float projection[16]; + float inverse_projection[16]; + float size[2]; + float inv_size[2]; + float depth_scale[2]; + float effect_radius; + float intensity; + float slice_count; + float samples_per_slice_side; + uint32_t debug_view; + float _pad; +}; +static_assert(sizeof(AoUniforms) % 16 == 0); + +struct ComputePayload { + WGPUTextureView depth; // frame-pooled scene depth snapshot + WGPUTextureView preprocessedDepthMips[5]; + WGPUTextureView preprocessedDepthAll; + WGPUTextureView aoNoisy; + WGPUTextureView depthDifferences; + WGPUTextureView aoFinal; + uint32_t uniform_offset; + uint32_t uniform_size; + uint32_t width; + uint32_t height; +}; +static_assert(sizeof(ComputePayload) <= GFX_INLINE_DRAW_PAYLOAD_SIZE); +static_assert(std::is_trivially_copyable_v); + +struct CompositePayload { + WGPUTextureView aoFinal; + WGPUTextureView preprocessedDepth; // debug views reconstruct normals/depth from it + WGPUTextureView sceneDepth; // raw snapshot, for the bypass debug views + uint32_t uniform_offset; + uint32_t uniform_size; + uint32_t debug_view; +}; +static_assert(sizeof(CompositePayload) <= GFX_INLINE_DRAW_PAYLOAD_SIZE); +static_assert(std::is_trivially_copyable_v); + +int64_t get_int_option(ConfigVarHandle handle, int64_t fallback) { + int64_t value = fallback; + if (handle == 0 || svc_config->get_int(mod_ctx, handle, &value) != MOD_OK) { + return fallback; + } + return value; +} + +bool get_bool_option(ConfigVarHandle handle, bool fallback) { + bool value = fallback; + if (handle == 0 || svc_config->get_bool(mod_ctx, handle, &value) != MOD_OK) { + return fallback; + } + return value; +} + +// XeGTAO/Bevy quality presets: slices x (samples per slice side * 2). +void quality_counts(int64_t quality, float& sliceCount, float& samplesPerSliceSide) { + switch (std::clamp(quality, 0, 3)) { + case 0: + sliceCount = 1.0f; + samplesPerSliceSide = 2.0f; + break; + case 1: + sliceCount = 2.0f; + samplesPerSliceSide = 2.0f; + break; + default: + case 2: + sliceCount = 3.0f; + samplesPerSliceSide = 3.0f; + break; + case 3: + sliceCount = 9.0f; + samplesPerSliceSide = 3.0f; + break; + } +} + +WGPUShaderModule create_shader_module(const char* label, const ResourceBuffer& source) { + WGPUShaderSourceWGSL wgsl = WGPU_SHADER_SOURCE_WGSL_INIT; + wgsl.code = {static_cast(source.data), source.size}; + WGPUShaderModuleDescriptor moduleDesc = WGPU_SHADER_MODULE_DESCRIPTOR_INIT; + moduleDesc.nextInChain = &wgsl.chain; + moduleDesc.label = {label, WGPU_STRLEN}; + return wgpuDeviceCreateShaderModule(g_deviceInfo.device, &moduleDesc); +} + +bool build_compute_pipeline(const char* label, const ResourceBuffer& source, const char* entry, + WGPUComputePipeline& outPipeline, WGPUBindGroupLayout& outLayout) { + WGPUShaderModule module = create_shader_module(label, source); + if (module == nullptr) { + return false; + } + WGPUComputePipelineDescriptor pipelineDesc = WGPU_COMPUTE_PIPELINE_DESCRIPTOR_INIT; + pipelineDesc.label = {label, WGPU_STRLEN}; + pipelineDesc.compute.module = module; + pipelineDesc.compute.entryPoint = {entry, WGPU_STRLEN}; + outPipeline = wgpuDeviceCreateComputePipeline(g_deviceInfo.device, &pipelineDesc); + wgpuShaderModuleRelease(module); + if (outPipeline == nullptr) { + return false; + } + outLayout = wgpuComputePipelineGetBindGroupLayout(outPipeline, 0); + return outLayout != nullptr; +} + +bool build_composite_pipeline( + bool blend, WGPURenderPipeline& outPipeline, WGPUBindGroupLayout& outLayout) { + WGPUShaderModule module = create_shader_module("AO composite", g_compositeSource); + if (module == nullptr) { + return false; + } + + // Multiply blend + WGPUBlendState blendState{ + .color = + { + .operation = WGPUBlendOperation_Add, + .srcFactor = WGPUBlendFactor_Dst, + .dstFactor = WGPUBlendFactor_Zero, + }, + .alpha = + { + .operation = WGPUBlendOperation_Add, + .srcFactor = WGPUBlendFactor_Zero, + .dstFactor = WGPUBlendFactor_One, + }, + }; + WGPUColorTargetState colorTarget = WGPU_COLOR_TARGET_STATE_INIT; + colorTarget.format = g_deviceInfo.color_format; + if (blend) { + colorTarget.blend = &blendState; + } + WGPUFragmentState fragment = WGPU_FRAGMENT_STATE_INIT; + fragment.module = module; + fragment.entryPoint = {"fs_main", WGPU_STRLEN}; + fragment.targetCount = 1; + fragment.targets = &colorTarget; + // Depth state must match the EFB pass despite never touching depth. + WGPUDepthStencilState depthStencil = WGPU_DEPTH_STENCIL_STATE_INIT; + depthStencil.format = g_deviceInfo.depth_format; + depthStencil.depthWriteEnabled = WGPUOptionalBool_False; + depthStencil.depthCompare = WGPUCompareFunction_Always; + + WGPURenderPipelineDescriptor pipelineDesc = WGPU_RENDER_PIPELINE_DESCRIPTOR_INIT; + pipelineDesc.label = {blend ? "AO composite" : "AO composite (debug)", WGPU_STRLEN}; + pipelineDesc.vertex.module = module; + pipelineDesc.vertex.entryPoint = {"vs_main", WGPU_STRLEN}; + pipelineDesc.primitive.topology = WGPUPrimitiveTopology_TriangleList; + pipelineDesc.depthStencil = &depthStencil; + pipelineDesc.multisample.count = g_deviceInfo.sample_count; + pipelineDesc.fragment = &fragment; + outPipeline = wgpuDeviceCreateRenderPipeline(g_deviceInfo.device, &pipelineDesc); + wgpuShaderModuleRelease(module); + if (outPipeline == nullptr) { + return false; + } + outLayout = wgpuRenderPipelineGetBindGroupLayout(outPipeline, 0); + return outLayout != nullptr; +} + +// Hilbert curve index LUT for the R2 noise sequence, generated once at init. +// Ported from Bevy's generate_hilbert_index_lut (https://www.shadertoy.com/view/3tB3z3). +uint16_t hilbert_index(uint16_t x, uint16_t y) { + uint16_t index = 0; + for (uint16_t level = 32; level > 0; level /= 2) { + const uint16_t regionX = (x & level) > 0 ? 1 : 0; + const uint16_t regionY = (y & level) > 0 ? 1 : 0; + index += level * level * ((3 * regionX) ^ regionY); + if (regionY == 0) { + if (regionX == 1) { + x = 63 - x; + y = 63 - y; + } + std::swap(x, y); + } + } + return index; +} + +bool build_hilbert_lut() { + WGPUTextureDescriptor texDesc = WGPU_TEXTURE_DESCRIPTOR_INIT; + texDesc.label = {"AO hilbert LUT", WGPU_STRLEN}; + texDesc.usage = WGPUTextureUsage_TextureBinding | WGPUTextureUsage_CopyDst; + texDesc.size = {64, 64, 1}; + texDesc.format = WGPUTextureFormat_R16Uint; + g_hilbertLut = wgpuDeviceCreateTexture(g_deviceInfo.device, &texDesc); + if (g_hilbertLut == nullptr) { + return false; + } + g_hilbertLutView = wgpuTextureCreateView(g_hilbertLut, nullptr); + if (g_hilbertLutView == nullptr) { + return false; + } + + uint16_t lut[64 * 64]; + for (uint16_t y = 0; y < 64; ++y) { + for (uint16_t x = 0; x < 64; ++x) { + lut[y * 64 + x] = hilbert_index(x, y); + } + } + WGPUTexelCopyTextureInfo dst = WGPU_TEXEL_COPY_TEXTURE_INFO_INIT; + dst.texture = g_hilbertLut; + WGPUTexelCopyBufferLayout layout{.offset = 0, .bytesPerRow = 64 * 2, .rowsPerImage = 64}; + WGPUExtent3D extent{64, 64, 1}; + wgpuQueueWriteTexture(g_deviceInfo.queue, &dst, lut, sizeof(lut), &layout, &extent); + return true; +} + +void release_targets(AoTargets& targets) { + for (auto*& view : targets.preprocessedDepthMips) { + if (view != nullptr) { + wgpuTextureViewRelease(view); + view = nullptr; + } + } + const auto releaseView = [](WGPUTextureView& view) { + if (view != nullptr) { + wgpuTextureViewRelease(view); + view = nullptr; + } + }; + const auto releaseTexture = [](WGPUTexture& texture) { + if (texture != nullptr) { + wgpuTextureRelease(texture); + texture = nullptr; + } + }; + releaseView(targets.preprocessedDepthAll); + releaseView(targets.aoNoisyView); + releaseView(targets.depthDifferencesView); + releaseView(targets.aoFinalView); + releaseTexture(targets.preprocessedDepth); + releaseTexture(targets.aoNoisy); + releaseTexture(targets.depthDifferences); + releaseTexture(targets.aoFinal); + targets.width = targets.height = 0; +} + +void tick_retired_targets() { + for (auto it = g_retiredTargets.begin(); it != g_retiredTargets.end();) { + if (--it->framesLeft <= 0) { + release_targets(it->targets); + it = g_retiredTargets.erase(it); + } else { + ++it; + } + } +} + +bool ensure_targets(uint32_t width, uint32_t height) { + if (g_targets.width == width && g_targets.height == height) { + return true; + } + if (g_targets.width != 0) { + g_retiredTargets.push_back(RetiredTargets{std::exchange(g_targets, AoTargets{}), 4}); + } + + const auto createStorageTexture = [&](const char* label, WGPUTextureFormat format, + uint32_t mipCount, WGPUTexture& outTexture) { + WGPUTextureDescriptor texDesc = WGPU_TEXTURE_DESCRIPTOR_INIT; + texDesc.label = {label, WGPU_STRLEN}; + texDesc.usage = WGPUTextureUsage_StorageBinding | WGPUTextureUsage_TextureBinding; + texDesc.size = {width, height, 1}; + texDesc.format = format; + texDesc.mipLevelCount = mipCount; + outTexture = wgpuDeviceCreateTexture(g_deviceInfo.device, &texDesc); + return outTexture != nullptr; + }; + + bool ok = createStorageTexture("AO preprocessed depth", WGPUTextureFormat_R32Float, 5, + g_targets.preprocessedDepth) && + createStorageTexture("AO noisy", WGPUTextureFormat_R32Float, 1, g_targets.aoNoisy) && + createStorageTexture("AO depth differences", WGPUTextureFormat_R32Uint, 1, + g_targets.depthDifferences) && + createStorageTexture("AO final", WGPUTextureFormat_R32Float, 1, g_targets.aoFinal); + if (ok) { + for (uint32_t mip = 0; mip < 5 && ok; ++mip) { + WGPUTextureViewDescriptor viewDesc = WGPU_TEXTURE_VIEW_DESCRIPTOR_INIT; + viewDesc.baseMipLevel = mip; + viewDesc.mipLevelCount = 1; + g_targets.preprocessedDepthMips[mip] = + wgpuTextureCreateView(g_targets.preprocessedDepth, &viewDesc); + ok = g_targets.preprocessedDepthMips[mip] != nullptr; + } + } + if (ok) { + g_targets.preprocessedDepthAll = + wgpuTextureCreateView(g_targets.preprocessedDepth, nullptr); + g_targets.aoNoisyView = wgpuTextureCreateView(g_targets.aoNoisy, nullptr); + g_targets.depthDifferencesView = wgpuTextureCreateView(g_targets.depthDifferences, nullptr); + g_targets.aoFinalView = wgpuTextureCreateView(g_targets.aoFinal, nullptr); + ok = g_targets.preprocessedDepthAll != nullptr && g_targets.aoNoisyView != nullptr && + g_targets.depthDifferencesView != nullptr && g_targets.aoFinalView != nullptr; + } + if (!ok) { + release_targets(g_targets); + return false; + } + g_targets.width = width; + g_targets.height = height; + return true; +} + +constexpr uint32_t div_ceil(uint32_t numerator, uint32_t denominator) { + return (numerator + denominator - 1) / denominator; +} + +// Render worker thread: the AO chain as one compute pass with three dispatches. +void on_compute( + ModContext*, const GfxComputeContext* ctx, const void* payload, size_t payloadSize, void*) { + if (payloadSize != sizeof(ComputePayload)) { + return; + } + ComputePayload data; + std::memcpy(&data, payload, sizeof(data)); + if (data.depth == nullptr || g_preprocessPipeline == nullptr) { + return; + } + + const auto makeBindGroup = [&](WGPUBindGroupLayout layout, + std::initializer_list entries) { + WGPUBindGroupDescriptor bindGroupDesc = WGPU_BIND_GROUP_DESCRIPTOR_INIT; + bindGroupDesc.layout = layout; + bindGroupDesc.entryCount = entries.size(); + bindGroupDesc.entries = entries.begin(); + return wgpuDeviceCreateBindGroup(ctx->device, &bindGroupDesc); + }; + const auto textureEntry = [](uint32_t binding, WGPUTextureView view) { + WGPUBindGroupEntry entry = WGPU_BIND_GROUP_ENTRY_INIT; + entry.binding = binding; + entry.textureView = view; + return entry; + }; + const auto uniformEntry = [&](uint32_t binding) { + WGPUBindGroupEntry entry = WGPU_BIND_GROUP_ENTRY_INIT; + entry.binding = binding; + entry.buffer = ctx->uniform_buffer; + entry.offset = data.uniform_offset; + entry.size = data.uniform_size; + return entry; + }; + + WGPUBindGroup preprocessGroup = makeBindGroup(g_preprocessLayout, + {textureEntry(0, data.depth), textureEntry(1, data.preprocessedDepthMips[0]), + textureEntry(2, data.preprocessedDepthMips[1]), + textureEntry(3, data.preprocessedDepthMips[2]), + textureEntry(4, data.preprocessedDepthMips[3]), uniformEntry(5)}); + WGPUBindGroup mip4Group = + makeBindGroup(g_mip4Layout, {textureEntry(6, data.preprocessedDepthMips[3]), + textureEntry(7, data.preprocessedDepthMips[4])}); + WGPUBindGroup gtaoGroup = makeBindGroup( + g_gtaoLayout, {textureEntry(0, data.preprocessedDepthAll), + textureEntry(1, g_hilbertLutView), textureEntry(2, data.aoNoisy), + textureEntry(3, data.depthDifferences), uniformEntry(4)}); + WGPUBindGroup denoiseGroup = makeBindGroup( + g_denoiseLayout, {textureEntry(0, data.aoNoisy), textureEntry(1, data.depthDifferences), + textureEntry(2, data.aoFinal), uniformEntry(3)}); + if (preprocessGroup == nullptr || mip4Group == nullptr || gtaoGroup == nullptr || + denoiseGroup == nullptr) + { + const auto release = [](WGPUBindGroup group) { + if (group != nullptr) { + wgpuBindGroupRelease(group); + } + }; + release(preprocessGroup); + release(mip4Group); + release(gtaoGroup); + release(denoiseGroup); + return; + } + + WGPUComputePassDescriptor passDesc = WGPU_COMPUTE_PASS_DESCRIPTOR_INIT; + passDesc.label = {"AO chain", WGPU_STRLEN}; + WGPUComputePassEncoder pass = wgpuCommandEncoderBeginComputePass(ctx->encoder, &passDesc); + // Each preprocess workgroup covers 16x16 MIP-0 texels (8x8 invocations, 2x2 texels each). + wgpuComputePassEncoderSetPipeline(pass, g_preprocessPipeline); + wgpuComputePassEncoderSetBindGroup(pass, 0, preprocessGroup, 0, nullptr); + wgpuComputePassEncoderDispatchWorkgroups( + pass, div_ceil(data.width, 16), div_ceil(data.height, 16), 1); + wgpuComputePassEncoderSetPipeline(pass, g_mip4Pipeline); + wgpuComputePassEncoderSetBindGroup(pass, 0, mip4Group, 0, nullptr); + wgpuComputePassEncoderDispatchWorkgroups(pass, div_ceil(std::max(data.width >> 4, 1u), 8), + div_ceil(std::max(data.height >> 4, 1u), 8), 1); + wgpuComputePassEncoderSetPipeline(pass, g_gtaoPipeline); + wgpuComputePassEncoderSetBindGroup(pass, 0, gtaoGroup, 0, nullptr); + wgpuComputePassEncoderDispatchWorkgroups( + pass, div_ceil(data.width, 8), div_ceil(data.height, 8), 1); + wgpuComputePassEncoderSetPipeline(pass, g_denoisePipeline); + wgpuComputePassEncoderSetBindGroup(pass, 0, denoiseGroup, 0, nullptr); + wgpuComputePassEncoderDispatchWorkgroups( + pass, div_ceil(data.width, 8), div_ceil(data.height, 8), 1); + wgpuComputePassEncoderEnd(pass); + wgpuComputePassEncoderRelease(pass); + + wgpuBindGroupRelease(preprocessGroup); + wgpuBindGroupRelease(mip4Group); + wgpuBindGroupRelease(gtaoGroup); + wgpuBindGroupRelease(denoiseGroup); + g_chainExecuted.store(true, std::memory_order_release); +} + +// Render worker thread: composite the AO over the scene (or show it, in debug view). +void on_draw( + ModContext*, const GfxDrawContext* ctx, const void* payload, size_t payloadSize, void*) { + if (payloadSize != sizeof(CompositePayload)) { + return; + } + CompositePayload data; + std::memcpy(&data, payload, sizeof(data)); + WGPURenderPipeline pipeline = + data.debug_view != 0 ? g_compositeDebugPipeline : g_compositePipeline; + WGPUBindGroupLayout layout = data.debug_view != 0 ? g_compositeDebugLayout : g_compositeLayout; + if (data.aoFinal == nullptr || data.preprocessedDepth == nullptr || + data.sceneDepth == nullptr || pipeline == nullptr) + { + return; + } + + WGPUBindGroupEntry entries[4] = {WGPU_BIND_GROUP_ENTRY_INIT, WGPU_BIND_GROUP_ENTRY_INIT, + WGPU_BIND_GROUP_ENTRY_INIT, WGPU_BIND_GROUP_ENTRY_INIT}; + entries[0].binding = 0; + entries[0].textureView = data.aoFinal; + entries[1].binding = 1; + entries[1].textureView = data.preprocessedDepth; + entries[2].binding = 2; + entries[2].textureView = data.sceneDepth; + entries[3].binding = 3; + entries[3].buffer = ctx->uniform_buffer; + entries[3].offset = data.uniform_offset; + entries[3].size = data.uniform_size; + WGPUBindGroupDescriptor bindGroupDesc = WGPU_BIND_GROUP_DESCRIPTOR_INIT; + bindGroupDesc.layout = layout; + bindGroupDesc.entryCount = 4; + bindGroupDesc.entries = entries; + WGPUBindGroup bindGroup = wgpuDeviceCreateBindGroup(ctx->device, &bindGroupDesc); + if (bindGroup == nullptr) { + return; + } + + wgpuRenderPassEncoderSetPipeline(ctx->pass, pipeline); + wgpuRenderPassEncoderSetBindGroup(ctx->pass, 0, bindGroup, 0, nullptr); + wgpuRenderPassEncoderDraw(ctx->pass, 3, 1, 0, 0); + wgpuBindGroupRelease(bindGroup); +} + +// Game thread, after opaque scene draws and before translucent/fog overlay lists. +void on_scene_after_opaque(ModContext*, const GfxStageContext* stageCtx, void*) { + tick_retired_targets(); + if (!get_bool_option(g_cvarEnabled, true)) { + return; + } + if (stageCtx == nullptr || stageCtx->struct_size < sizeof(GfxStageContext) || + stageCtx->game_view == nullptr) + { + return; + } + + CameraInfo camera = CAMERA_INFO_INIT; + if (svc_camera->get_camera(mod_ctx, stageCtx->game_view, &camera) != MOD_OK) { + return; + } + + GfxResolveDesc resolveDesc = GFX_RESOLVE_DESC_INIT; + resolveDesc.color = false; + resolveDesc.depth = true; + GfxResolvedTargets resolved = GFX_RESOLVED_TARGETS_INIT; + if (svc_gfx->resolve_pass(mod_ctx, &resolveDesc, &resolved) != MOD_OK || + resolved.depth == nullptr) + { + if (!g_warnedNoDepth) { + g_warnedNoDepth = true; + svc_log->warn(mod_ctx, "depth snapshots unavailable; AO disabled"); + } + return; + } + + const bool halfRes = get_bool_option(g_cvarHalfRes, false); + const uint32_t divisor = halfRes ? 2 : 1; + const uint32_t width = resolved.width / divisor; + const uint32_t height = resolved.height / divisor; + if (width < 32 || height < 32 || !ensure_targets(width, height)) { + return; + } + + AoUniforms uniforms{}; + std::memcpy(uniforms.projection, camera.proj_from_view, sizeof(uniforms.projection)); + std::memcpy( + uniforms.inverse_projection, camera.view_from_proj, sizeof(uniforms.inverse_projection)); + uniforms.size[0] = static_cast(width); + uniforms.size[1] = static_cast(height); + uniforms.inv_size[0] = 1.0f / uniforms.size[0]; + uniforms.inv_size[1] = 1.0f / uniforms.size[1]; + uniforms.depth_scale[0] = static_cast(resolved.width) / uniforms.size[0]; + uniforms.depth_scale[1] = static_cast(resolved.height) / uniforms.size[1]; + uniforms.effect_radius = + static_cast(std::clamp(get_int_option(g_cvarRadius, 70), 10, 500)); + uniforms.intensity = + static_cast(std::clamp(get_int_option(g_cvarIntensity, 100), 0, 100)) / + 100.0f; + quality_counts( + get_int_option(g_cvarQuality, 2), uniforms.slice_count, uniforms.samples_per_slice_side); + const uint32_t debugMode = + static_cast(std::clamp(get_int_option(g_cvarDebugView, 0), 0, 4)); + uniforms.debug_view = debugMode; + + GfxRange uniformRange{0, 0}; + if (svc_gfx->push_uniform(mod_ctx, &uniforms, sizeof(uniforms), &uniformRange) != MOD_OK) { + return; + } + + ComputePayload computePayload{}; + computePayload.depth = resolved.depth; + for (int mip = 0; mip < 5; ++mip) { + computePayload.preprocessedDepthMips[mip] = g_targets.preprocessedDepthMips[mip]; + } + computePayload.preprocessedDepthAll = g_targets.preprocessedDepthAll; + computePayload.aoNoisy = g_targets.aoNoisyView; + computePayload.depthDifferences = g_targets.depthDifferencesView; + computePayload.aoFinal = g_targets.aoFinalView; + computePayload.uniform_offset = uniformRange.offset; + computePayload.uniform_size = uniformRange.size; + computePayload.width = width; + computePayload.height = height; + if (svc_gfx->push_compute(mod_ctx, g_computeType, &computePayload, sizeof(computePayload)) != + MOD_OK) + { + return; + } + + const CompositePayload drawPayload{g_targets.aoFinalView, g_targets.preprocessedDepthAll, + resolved.depth, uniformRange.offset, uniformRange.size, debugMode}; + svc_gfx->push_draw(mod_ctx, g_drawType, &drawPayload, sizeof(drawPayload)); +} + +void add_control(UiElementHandle pane, const UiControlDesc& desc) { + svc_ui->pane_add_control(mod_ctx, pane, &desc, nullptr); +} + +void add_toggle(UiElementHandle pane, const char* label, ConfigVarHandle cvar, const char* help) { + UiControlDesc control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_TOGGLE; + control.label = label; + control.help_rml = help; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = cvar; + add_control(pane, control); +} + +ModResult build_controls_tab( + ModContext*, UiWindowHandle, UiElementHandle left, UiElementHandle right, void*, ModError*) { + (void)right; + + svc_ui->pane_add_section(mod_ctx, left, "Ambient Occlusion"); + add_toggle(left, "Enabled", g_cvarEnabled, "Enables the GTAO pass."); + + static const char* kQualityOptions[] = {"Low", "Medium", "High", "Ultra"}; + UiControlDesc control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_SELECT; + control.label = "Quality"; + control.help_rml = "Horizon slices and samples per pixel (XeGTAO presets: 4/8/18/54 spp)."; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = g_cvarQuality; + control.options = kQualityOptions; + control.option_count = 4; + add_control(left, control); + + control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_NUMBER; + control.label = "Radius"; + control.help_rml = "Occlusion sampling radius in world units."; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = g_cvarRadius; + control.min = 10; + control.max = 500; + control.step = 10; + add_control(left, control); + + control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_NUMBER; + control.label = "Intensity"; + control.help_rml = "How strongly occlusion darkens the scene."; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = g_cvarIntensity; + control.min = 0; + control.max = 100; + control.step = 5; + control.suffix = "%"; + add_control(left, control); + + add_toggle(left, "Half Resolution", g_cvarHalfRes, + "Computes AO at half resolution and upscales; faster, slightly softer."); + + static const char* kDebugOptions[] = {"Off", "AO", "Normals", "Depth", "Staircase"}; + control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_SELECT; + control.label = "Debug View"; + control.help_rml = "AO: raw visibility as grayscale.
Normals: the view-space " + "normals the GTAO pass consumes.
Depth: the preprocessed depth " + "as a distance gradient.
Staircase: detects quantized depth - smooth " + "depth is near-black with thin triangle edges, quantized depth lights " + "up across surfaces."; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = g_cvarDebugView; + control.options = kDebugOptions; + control.option_count = 5; + add_control(left, control); + return MOD_OK; +} + +void on_controls_window_closed(ModContext*, UiWindowHandle, void*) { + g_controlsWindow = 0; +} + +void on_open_controls(ModContext*, void*) { + if (g_controlsWindow != 0) { + return; + } + UiTabDesc tabs[1] = {UI_TAB_DESC_INIT}; + tabs[0].title = "Controls"; + tabs[0].build = build_controls_tab; + UiWindowDesc desc = UI_WINDOW_DESC_INIT; + desc.tabs = tabs; + desc.tab_count = 1; + desc.on_closed = on_controls_window_closed; + if (svc_ui->window_push(mod_ctx, &desc, &g_controlsWindow) != MOD_OK) { + svc_log->error(mod_ctx, "failed to open AO controls window"); + } +} + +ModResult build_panel(ModContext*, UiElementHandle panel, void*, ModError*) { + UiControlDesc control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_TOGGLE; + control.label = "Enabled"; + control.binding = UI_BINDING_CONFIG_VAR; + control.config_var = g_cvarEnabled; + add_control(panel, control); + + control = UI_CONTROL_DESC_INIT; + control.kind = UI_CONTROL_BUTTON; + control.label = "Open Controls"; + control.on_pressed = on_open_controls; + add_control(panel, control); + return MOD_OK; +} + +ModResult register_bool_option( + const char* name, bool defaultValue, ConfigVarHandle& outHandle, ModError* error) { + ConfigVarDesc cvarDesc = CONFIG_VAR_DESC_INIT; + cvarDesc.name = name; + cvarDesc.type = CONFIG_VAR_BOOL; + cvarDesc.default_bool = defaultValue; + if (svc_config->register_var(mod_ctx, &cvarDesc, &outHandle) != MOD_OK) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to register AO option"); + } + return MOD_OK; +} + +ModResult register_int_option( + const char* name, int64_t defaultValue, ConfigVarHandle& outHandle, ModError* error) { + ConfigVarDesc cvarDesc = CONFIG_VAR_DESC_INIT; + cvarDesc.name = name; + cvarDesc.type = CONFIG_VAR_INT; + cvarDesc.default_int = defaultValue; + if (svc_config->register_var(mod_ctx, &cvarDesc, &outHandle) != MOD_OK) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to register AO option"); + } + return MOD_OK; +} + +} // namespace + +extern "C" { + +MOD_EXPORT ModResult mod_initialize(ModError* error) { + ModResult result = svc_resource->load(mod_ctx, "preprocess_depth.wgsl", &g_preprocessSource); + if (result == MOD_OK) { + result = svc_resource->load(mod_ctx, "gtao.wgsl", &g_gtaoSource); + } + if (result == MOD_OK) { + result = svc_resource->load(mod_ctx, "denoise.wgsl", &g_denoiseSource); + } + if (result == MOD_OK) { + result = svc_resource->load(mod_ctx, "composite.wgsl", &g_compositeSource); + } + if (result != MOD_OK) { + return dusk::mods::set_error(error, result, "failed to load AO shaders"); + } + + result = register_bool_option("effectEnabled", false, g_cvarEnabled, error); + if (result != MOD_OK) { + return result; + } + result = register_int_option("quality", 2, g_cvarQuality, error); + if (result != MOD_OK) { + return result; + } + result = register_int_option("radius", 70, g_cvarRadius, error); + if (result != MOD_OK) { + return result; + } + result = register_int_option("intensity", 100, g_cvarIntensity, error); + if (result != MOD_OK) { + return result; + } + result = register_bool_option("halfRes", false, g_cvarHalfRes, error); + if (result != MOD_OK) { + return result; + } + result = register_int_option("debugMode", 0, g_cvarDebugView, error); + if (result != MOD_OK) { + return result; + } + + if (svc_gfx->get_device_info(mod_ctx, &g_deviceInfo) != MOD_OK) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to query device info"); + } + if (!build_compute_pipeline("AO preprocess depth", g_preprocessSource, "preprocess_depth", + g_preprocessPipeline, g_preprocessLayout) || + !build_compute_pipeline("AO downsample mip4", g_preprocessSource, "downsample_mip4", + g_mip4Pipeline, g_mip4Layout) || + !build_compute_pipeline("AO gtao", g_gtaoSource, "gtao", g_gtaoPipeline, g_gtaoLayout) || + !build_compute_pipeline( + "AO denoise", g_denoiseSource, "spatial_denoise", g_denoisePipeline, g_denoiseLayout)) + { + return dusk::mods::set_error(error, MOD_ERROR, "failed to create AO compute pipelines"); + } + if (!build_composite_pipeline(true, g_compositePipeline, g_compositeLayout) || + !build_composite_pipeline(false, g_compositeDebugPipeline, g_compositeDebugLayout)) + { + return dusk::mods::set_error(error, MOD_ERROR, "failed to create AO composite pipeline"); + } + if (!build_hilbert_lut()) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to create AO noise LUT"); + } + + GfxComputeTypeDesc computeDesc = GFX_COMPUTE_TYPE_DESC_INIT; + computeDesc.label = "AO chain"; + computeDesc.callback = on_compute; + if (svc_gfx->register_compute_type(mod_ctx, &computeDesc, &g_computeType) != MOD_OK) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to register compute type"); + } + GfxDrawTypeDesc drawDesc = GFX_DRAW_TYPE_DESC_INIT; + drawDesc.label = "AO composite"; + drawDesc.draw = on_draw; + if (svc_gfx->register_draw_type(mod_ctx, &drawDesc, &g_drawType) != MOD_OK) { + return dusk::mods::set_error(error, MOD_ERROR, "failed to register draw type"); + } + GfxStageHookDesc stageDesc = GFX_STAGE_HOOK_DESC_INIT; + stageDesc.callback = on_scene_after_opaque; + if (svc_gfx->register_stage_hook( + mod_ctx, GFX_STAGE_SCENE_AFTER_OPAQUE, &stageDesc, &g_afterOpaqueHook) != MOD_OK) + { + return dusk::mods::set_error(error, MOD_ERROR, "failed to register stage hook"); + } + + UiModsPanelDesc panelDesc = UI_MODS_PANEL_DESC_INIT; + panelDesc.build = build_panel; + svc_ui->register_mods_panel(mod_ctx, &panelDesc); + + svc_log->info(mod_ctx, "ao_mod ready"); + return MOD_OK; +} + +MOD_EXPORT ModResult mod_update(ModError*) { + if (!g_loggedChain && g_chainExecuted.load(std::memory_order_acquire)) { + g_loggedChain = true; + svc_log->info(mod_ctx, "AO chain executed OK"); + } + return MOD_OK; +} + +MOD_EXPORT ModResult mod_shutdown(ModError*) { + svc_resource->free(mod_ctx, &g_preprocessSource); + svc_resource->free(mod_ctx, &g_gtaoSource); + svc_resource->free(mod_ctx, &g_denoiseSource); + svc_resource->free(mod_ctx, &g_compositeSource); + + release_targets(g_targets); + for (auto& retired : g_retiredTargets) { + release_targets(retired.targets); + } + g_retiredTargets.clear(); + + const auto releasePipeline = [](WGPUComputePipeline& pipeline) { + if (pipeline != nullptr) { + wgpuComputePipelineRelease(pipeline); + pipeline = nullptr; + } + }; + const auto releaseLayout = [](WGPUBindGroupLayout& layout) { + if (layout != nullptr) { + wgpuBindGroupLayoutRelease(layout); + layout = nullptr; + } + }; + releasePipeline(g_preprocessPipeline); + releasePipeline(g_mip4Pipeline); + releasePipeline(g_gtaoPipeline); + releasePipeline(g_denoisePipeline); + releaseLayout(g_preprocessLayout); + releaseLayout(g_mip4Layout); + releaseLayout(g_gtaoLayout); + releaseLayout(g_denoiseLayout); + if (g_compositePipeline != nullptr) { + wgpuRenderPipelineRelease(g_compositePipeline); + g_compositePipeline = nullptr; + } + if (g_compositeDebugPipeline != nullptr) { + wgpuRenderPipelineRelease(g_compositeDebugPipeline); + g_compositeDebugPipeline = nullptr; + } + releaseLayout(g_compositeLayout); + releaseLayout(g_compositeDebugLayout); + if (g_hilbertLutView != nullptr) { + wgpuTextureViewRelease(g_hilbertLutView); + g_hilbertLutView = nullptr; + } + if (g_hilbertLut != nullptr) { + wgpuTextureRelease(g_hilbertLut); + g_hilbertLut = nullptr; + } + g_cvarEnabled = g_cvarQuality = g_cvarRadius = g_cvarIntensity = 0; + g_cvarHalfRes = g_cvarDebugView = 0; + g_computeType = g_drawType = 0; + g_afterOpaqueHook = 0; + g_controlsWindow = 0; + return MOD_OK; +} +} diff --git a/tools/mod_template/CMakeLists.txt b/mods/template_mod/CMakeLists.txt similarity index 100% rename from tools/mod_template/CMakeLists.txt rename to mods/template_mod/CMakeLists.txt diff --git a/mods/template_mod/mod.json b/mods/template_mod/mod.json new file mode 100644 index 0000000000..e60714edf5 --- /dev/null +++ b/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" +} diff --git a/tools/mod_template/res/.gitkeep b/mods/template_mod/res/.gitkeep similarity index 100% rename from tools/mod_template/res/.gitkeep rename to mods/template_mod/res/.gitkeep diff --git a/tools/mod_template/src/mod.cpp b/mods/template_mod/src/mod.cpp similarity index 100% rename from tools/mod_template/src/mod.cpp rename to mods/template_mod/src/mod.cpp diff --git a/sdk/CMakeLists.txt b/sdk/CMakeLists.txt index 533bcc22c1..8670603866 100644 --- a/sdk/CMakeLists.txt +++ b/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(/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") diff --git a/tools/mod_template/mod.json b/tools/mod_template/mod.json deleted file mode 100644 index 46507a87ee..0000000000 --- a/tools/mod_template/mod.json +++ /dev/null @@ -1,7 +0,0 @@ -{ - "id": "com.example.mod", - "name": "Template Mod", - "version": "1.0.0", - "author": "You", - "description": "An example Dusklight mod" -}