From 9ba1fd460b5e814f17b8c719bb357d3c6468421f Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 15:58:15 +0200 Subject: [PATCH 01/16] feat(pt): mode plumbing + roadmap for the three-tier path tracer MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Groundwork for docs/pt/pt-roadmap.md — no rendering yet, just the switch: - renderer.pt_mode (0 off / 1 progressive / 2 realtime), set_path_tracing, pt_supported (== hw_rt_enabled: same ray-query + TLAS requirement as Lumen's HW trace), pt_active. Deliberately no software fallback: a software path tracer is seconds per frame, worse than useless in any mode we would ship. - FFI bloom_set_path_tracing / bloom_path_tracing_supported (+ manifest). - TS setPathTracing / isPathTracingSupported in bloom/core. The roadmap fixes the design: one compute megakernel (wgpu 29 has inline ray query only — no RT pipelines, no SBT), primary hits from the G-buffer, Tier 2 via a concatenated geometry megabuffer rather than buffer binding arrays, Tier 3 capped at temporal+à-trous (no ML denoiser exists for wgpu), and bloom-reference as the acceptance oracle with its BRDF conventions inherited so images stay comparable number-to-number. --- docs/pt/pt-roadmap.md | 128 +++++++++++++++++++++++++++ native/shared/src/ffi_core/visual.rs | 20 +++++ native/shared/src/renderer/mod.rs | 29 ++++++ package.json | 12 +++ src/core/index.ts | 21 +++++ 5 files changed, 210 insertions(+) create mode 100644 docs/pt/pt-roadmap.md diff --git a/docs/pt/pt-roadmap.md b/docs/pt/pt-roadmap.md new file mode 100644 index 0000000..b42629d --- /dev/null +++ b/docs/pt/pt-roadmap.md @@ -0,0 +1,128 @@ +# Path tracing — three tiers, one kernel + +A real GPU path tracer for Bloom, built on the ray-query/TLAS infrastructure +that Lumen already maintains. Three tiers share one megakernel; they differ in +what a ray hit can know (Tier 2) and in how few samples we can get away with +per frame (Tier 3). + +This is **in addition to** Lumen, never instead of it. Lumen remains the +default gameplay GI on every platform; path tracing is a render *mode* that +engages where the hardware makes it honest, and falls back cleanly where it +doesn't. + +## Why this shape + +wgpu 29 exposes **inline ray query only** — no ray-tracing pipelines, no +shader binding tables, no `traceRay`, no hit shaders. So the tracer is a +compute megakernel that owns its own bounce loop and material dispatch. That +is a constraint, but it is also the portable shape: the same WGSL runs on +DXR 1.1, Vulkan ray query, and Metal. + +The CPU reference tracer (`tools/bloom-reference`) is the acceptance oracle +for all of this. Two independent implementations agreeing is the strongest +correctness signal available to us — and the GPU tracer inherits the +reference's conventions (GGX + Burley, NEE with MIS balance heuristic, ACES + +fixed exposure) so images are comparable number-to-number. + +## Modes + +| Mode | Who | What | +|---|---|---| +| `0` off | default | Lumen GI as today | +| `1` progressive | editor "final quality" view, stills, ground truth in-engine | 1 spp/frame accumulated indefinitely while the camera is still; reset on move/scene change. No denoiser — convergence is the denoiser. | +| `2` realtime | gameplay on RT-capable GPUs | 1 spp/frame, temporal reprojection accumulation + à-trous spatial filter; optional half-res. Honest label: denoised 1-spp path tracing, the same family as Quake II RTX. | + +Mode is a runtime setting (`bloom_set_path_tracing`), toggleable per frame. +When PT is active, SSGI/SSR/GTAO are skipped (their output would be +overwritten; skipping banks their cost). Shadow cascades still render — mode 2 +reuses them nowhere, but the card-light pass (Tier 1 hit shading) does. + +## Tiers + +### Tier 1 — PT-1: progressive megakernel (correct diffuse transport) + +Primary hit reconstructed from the **G-buffer** (free, and sharper than traced +primaries). Then per pixel per frame: + +- NEE to the **sun**: sample the solar disc (~0.53°), one `rayQuery` shadow + ray. Sky misses exclude the disc so BSDF samples cannot double-count it. +- NEE to **point lights**: sample one light per bounce, inverse-square + + range falloff matched to the raster path. +- **Diffuse bounce**: cosine-weighted hemisphere; hit shading from the Mesh + Card albedo atlas (textured) or flat instance albedo — the same + `InstanceGiData` the Lumen HW trace reads. Emissive picked up at hits. +- Russian roulette from bounce 3; max 8 bounces (progressive) / 2 (realtime). +- Accumulate into rgba32float; reset when the view matrix moves beyond + epsilon or `tlas_version` bumps. + +Known honest limits at this tier: flat per-mesh hit normals and +card-resolution albedo → no sharp specular at secondary hits, no glass. +Correct, converged **diffuse** transport with real shadows and real +multi-bounce colour bleed. + +### Tier 2 — PT-2: real materials at the hit + +- **Geometry megabuffer**: at BLAS build, mesh positions/normals/UVs are also + appended to one shared storage buffer; instances carry + `{vtx_base, idx_base, material_id}`. One buffer — deliberately NOT + `binding_array` — so geometry fetch works wherever ray query does. +- Hit shading fetches the triangle, interpolates normal + UV. +- **Albedo textures** through `binding_array>`, gated on + `TEXTURE_BINDING_ARRAY` (+ non-uniform indexing); falls back to card + albedo where absent. DX12 with DXC and Vulkan both qualify. +- **GGX specular lobe** with lobe selection + MIS, mirroring + `bloom-reference/src/tracer.rs` so the two tracers stay comparable. +- Emissive from material data (VFX/muzzle flashes become real light in PT). + +### Tier 3 — PT-3/PT-4: gameplay + +- **PT-3**: temporal reprojection accumulation using the TAA velocity buffer + (disocclusion via depth/normal test, neighbourhood clamp, capped alpha), + then 3 à-trous iterations guided by depth + normal, firefly clamp. + Optional half-res trace + bilateral upsample (SSGI resolve pattern). +- **PT-4**: ReSTIR DI — per-pixel weighted reservoirs over light candidates, + temporal reuse with M-cap, spatial reuse over 3–5 neighbours, one shadow + ray for the winner. **Experimental flag until validated against + bloom-reference.** With today's ≤6 analytic lights plain NEE is nearly as + good; the payoff arrives when emissive particles/muzzle flashes become + lights (Tier 2). Do not oversell it before then. + +An ML denoiser is explicitly out of scope: nothing in wgpu runs vendor +denoisers, and a hand-rolled network is not happening. À-trous + temporal is +the honest ceiling here — it will look like early-RTX-era PT, not Cyberpunk +overdrive. That is the "as far as reasonable" line. + +## Platform / fallback matrix + +| Situation | Behaviour | +|---|---| +| `ray_query=true` (DX12+DXC, Vulkan RQ, Metal) | all modes available | +| `ray_query=false` | `bloom_set_path_tracing` is a no-op → Lumen; boot line + `bloom_pt_supported()` say why | +| `TEXTURE_BINDING_ARRAY` absent | Tier 2 textures fall back to card albedo; everything else unaffected | +| Web | never (no WebGPU RT); permanently Lumen SW | + +## Verification protocol (per the hard lessons) + +- **Binary-colour probes first**: every new pass proves it executes by + writing an unmistakable solid colour before any real math ships. No + image-diff conclusions under auto-exposure — PT comparisons run with + manual exposure, same as reference parity. +- **Convergence check**: progressive mode screenshot at ~2 s vs ~15 s; the + later one must be visibly smoother, not merely different. +- **Ground truth**: match camera + fixed exposure, render the same scene in + `bloom-reference`, compare with `bloom-diff` (RMSE/SSIM). Acceptance for + Tier 1: "close, differences explainable by flat hit normals / card-res + albedo." Tier 2 tightens that. +- **Perf on a fight, not the title screen** (perf round-3 lesson), and any + perf win larger than the change justifies = deleted geometry — screenshot + 3×. + +## Tickets + +| Ticket | Scope | +|---|---| +| PT-1 | Progressive megakernel, accumulation, FFI mode plumbing, shooter toggle | +| PT-2 | Geometry megabuffer, interpolated hit attributes, texture binding array, GGX+MIS, emissive hits | +| PT-3 | Realtime mode: temporal reprojection + à-trous, half-res option, perf gate | +| PT-4 | ReSTIR DI (experimental flag) | +| PT-5 | Settings/editor/gameplay integration, fallback matrix, reference-diff CI hook | diff --git a/native/shared/src/ffi_core/visual.rs b/native/shared/src/ffi_core/visual.rs index 264fb18..bbde564 100644 --- a/native/shared/src/ffi_core/visual.rs +++ b/native/shared/src/ffi_core/visual.rs @@ -340,6 +340,26 @@ macro_rules! __bloom_ffi_visual { }) } + // bloom_set_path_tracing — 0 off / 1 progressive / 2 realtime + // (docs/pt/pt-roadmap.md). Needs hardware ray query; without it + // the request is stored but nothing engages — check + // bloom_path_tracing_supported to know which world you are in. + #[no_mangle] + pub extern "C" fn bloom_set_path_tracing(mode: f64) { + $crate::ffi::guard("bloom_set_path_tracing", move || { + engine().renderer.set_path_tracing(mode as u32); + }) + } + + // bloom_path_tracing_supported — 1.0 when the device can trace + // (same ray-query requirement as Lumen's HW backend), else 0.0. + #[no_mangle] + pub extern "C" fn bloom_path_tracing_supported() -> f64 { + $crate::ffi::guard("bloom_path_tracing_supported", move || { + if engine().renderer.pt_supported() { 1.0 } else { 0.0 } + }) + } + // bloom_set_ssgi_intensity [source: macos] #[no_mangle] pub extern "C" fn bloom_set_ssgi_intensity(intensity: f64) { diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index a59c502..4488c08 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -1138,6 +1138,13 @@ pub struct Renderer { pub ssgi_radius: f32, /// SSGI master switch. pub ssgi_enabled: bool, + /// Path-tracing mode (docs/pt/pt-roadmap.md): 0 = off (Lumen as + /// usual), 1 = progressive (accumulate while the camera is still), + /// 2 = realtime (temporal + spatial denoise). Stores the *requested* + /// mode even without hardware ray query, so callers can query why + /// nothing engaged; the frame only switches when `pt_active()` says + /// both mode and hardware agree. + pub pt_mode: u32, /// EN-023 — mean flat albedo over the card instances; feeds the SW /// WSRC bake's ground-bounce term. Neutral mid-gray until the first /// instance-data upload computes the real scene average. @@ -6470,6 +6477,7 @@ impl Renderer { ssgi_intensity: 1.0, ssgi_radius: 20.0, ssgi_enabled: true, + pt_mode: 0, ssgi_backend_logged: None, gi_scene_avg_albedo: [0.35, 0.35, 0.35], probe_grid_w, @@ -7235,6 +7243,27 @@ impl Renderer { self.ssgi_enabled = on; } + /// Path-tracing mode request (0 off / 1 progressive / 2 realtime). + /// Clamped; anything above realtime means realtime. + pub fn set_path_tracing(&mut self, mode: u32) { + self.pt_mode = mode.min(2); + } + + /// Whether path tracing can run at all on this device: it needs the + /// same hardware ray query + TLAS the Lumen HW trace uses. + pub fn pt_supported(&self) -> bool { + self.hw_rt_enabled + } + + /// Whether the current frame should path-trace instead of running + /// the raster GI stack: a non-zero requested mode AND the hardware + /// to honour it. There is deliberately no software fallback — a + /// software path tracer would be seconds per frame, which is worse + /// than useless in every mode we ship. + pub fn pt_active(&self) -> bool { + self.pt_mode > 0 && self.pt_supported() + } + /// Ticket 013 — rasterise every pending mesh card into its /// assigned atlas slot. Called once per frame before the HW probe /// trace samples the atlas. Drains `scene.pending_card_captures`; diff --git a/package.json b/package.json index 36663d5..e194817 100644 --- a/package.json +++ b/package.json @@ -1662,6 +1662,18 @@ ], "returns": "void" }, + { + "name": "bloom_set_path_tracing", + "params": [ + "f64" + ], + "returns": "void" + }, + { + "name": "bloom_path_tracing_supported", + "params": [], + "returns": "f64" + }, { "name": "bloom_set_ssgi_intensity", "params": [ diff --git a/src/core/index.ts b/src/core/index.ts index f31ea58..fca58d1 100644 --- a/src/core/index.ts +++ b/src/core/index.ts @@ -39,6 +39,8 @@ declare function bloom_set_auto_resolution(targetHz: number, enabled: number): v declare function bloom_set_manual_exposure(value: number): void; declare function bloom_set_env_intensity(intensity: number): void; declare function bloom_set_ssgi_enabled(on: number): void; +declare function bloom_set_path_tracing(mode: number): void; +declare function bloom_path_tracing_supported(): number; declare function bloom_set_ssgi_intensity(intensity: number): void; declare function bloom_set_ssgi_radius(radius: number): void; declare function bloom_set_dof(enabled: number, focusDistance: number, aperture: number): void; @@ -409,6 +411,25 @@ export function setSsgiEnabled(on: boolean): void { bloom_set_ssgi_enabled(on ? 1 : 0); } +/** + * Path-tracing mode (engine docs/pt/pt-roadmap.md): + * 0 — off: the normal raster + Lumen pipeline (default). + * 1 — progressive: 1 sample/frame accumulated while the camera is still; + * resets on movement. Converges to ground truth — the "final quality" + * view for the editor and for stills. + * 2 — realtime: denoised 1-sample path tracing for gameplay. + * Requires hardware ray query (see isPathTracingSupported); on devices + * without it the request is a no-op and the engine stays on Lumen. + */ +export function setPathTracing(mode: number): void { + bloom_set_path_tracing(mode); +} + +/** True when the device can hardware-path-trace (ray query + TLAS). */ +export function isPathTracingSupported(): boolean { + return bloom_path_tracing_supported() !== 0; +} + /** SSGI intensity multiplier. 0 = off, 0.5 = default, 1+ = strong. */ export function setSsgiIntensity(intensity: number): void { bloom_set_ssgi_intensity(intensity); From b193e52b3c6d2b9dda792ef5667afd878b2025a8 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 16:34:08 +0200 Subject: [PATCH 02/16] feat(pt): PT-1 progressive path-trace megakernel (M1+M2 verified) One compute megakernel (shaders/pt.rs) replaces the lit opaque scene colour when path tracing is active: G-buffer primaries, TLAS bounce + NEE shadow rays, card-albedo hit shading, progressive accumulation (mode 1) or EMA (mode 2). Sky pixels are never written and the pass sits before translucency, so raster sky/clouds and water composite exactly as before. SSGI/SSR skip their dispatches while PT owns the frame; compose routes SSR to the cleared RT. Accumulation resets compare an UNJITTERED view-projection - the tracked current_vp_matrix carries the TAA Halton nudge, which reads as camera motion every frame and pins progressive mode at 1 sample. BLOOM_PT / BLOOM_PT_DEBUG env seed mode + debug view for batch runs. Verified on AMD 760M (DX12+DXC): magenta probe, standard-Z depth, convergence between 17s and 50s stills, live F9 off-switch restores raster bit-clean at 46 fps. Ticket doc carries the full status. --- docs/pt/PT-1-progressive-megakernel.md | 92 +++++ native/shared/src/renderer/formats.rs | 6 +- native/shared/src/renderer/mod.rs | 205 +++++++++- native/shared/src/renderer/postfx_chain.rs | 5 +- native/shared/src/renderer/pt_pass.rs | 175 +++++++++ native/shared/src/renderer/shaders/mod.rs | 2 + native/shared/src/renderer/shaders/pt.rs | 421 +++++++++++++++++++++ native/shared/src/renderer/ssgi_pass.rs | 6 +- native/shared/src/renderer/ssr_pass.rs | 7 +- 9 files changed, 912 insertions(+), 7 deletions(-) create mode 100644 docs/pt/PT-1-progressive-megakernel.md create mode 100644 native/shared/src/renderer/pt_pass.rs create mode 100644 native/shared/src/renderer/shaders/pt.rs diff --git a/docs/pt/PT-1-progressive-megakernel.md b/docs/pt/PT-1-progressive-megakernel.md new file mode 100644 index 0000000..13f03ae --- /dev/null +++ b/docs/pt/PT-1-progressive-megakernel.md @@ -0,0 +1,92 @@ +# PT-1 — Progressive path-trace megakernel (Tier 1) + +Parent: [pt-roadmap.md](pt-roadmap.md). Requires: hardware ray query (merged +in #91), TLAS + `InstanceGiData` + Mesh Card atlases (Lumen tickets 007b/013). + +## Scope + +One compute megakernel (`renderer/shaders/pt.rs`, pass in +`renderer/pt_pass.rs`) that replaces the lit scene colour when +`pt_active() && pt_mode == 1`: + +1. Primary hit from the G-buffer (depth → world pos, normal, albedo, + rough/metal). Sky pixels (far depth) get sky radiance directly. +2. Bounce loop, max 8: + - **NEE sun**: sample the solar disc (half-angle 0.265°), one shadow ray. + Delta-light radiometry identical to the raster sun so brightness + matches. Sky misses exclude the disc (no double count). + - **NEE point lights**: pick one of the frame's lights uniformly, + shadow ray, contribution / pdf. Same falloff as raster. + - **Diffuse bounce**: cosine hemisphere. Hit shading = card-atlas albedo + (textured) or flat `InstanceGiData` albedo + emissive — exactly what + the Lumen HW probe trace reads at hits today. + - Russian roulette from bounce 3 on throughput luminance. +3. Accumulate into an rgba32float buffer + (`accum = (accum * n + sample) / (n + 1)`); write tonemappable HDR into + scene colour. Reset `n` when the view matrix moves beyond epsilon or + `SceneGraph::tlas_version` bumps. + +While active: skip SSGI, SSR, GTAO dispatches (banked cost, and their +output would be overwritten). Shadow cascades keep rendering — the card +light pass needs them, and mode 0 must be switchable back per frame. + +## Milestone gates (each = compile + screenshot) + +- **M1**: kernel writes solid magenta gated on `pt_mode >= 1`. Proves + pipeline, bind groups, TLAS binding, dispatch, scene-colour write, F9. +- **M2**: sky + sun NEE + 1 diffuse bounce, accumulation working — 2 s vs + 15 s screenshots must show convergence (less noise, not merely + different). +- **M3**: full scope above. Point lights visible at night-ish exposure; + colour bleed visible at the building base (the EN-023 acceptance + criterion, finally measurable). + +## Acceptance + +- Convergence: variance visibly ↓ between 2 s and 15 s stills. +- Parity: same camera + fixed exposure in `bloom-reference` + (`--spp 256 --bounces 8`) vs converged in-engine frame; `bloom-diff` + RMSE reported in the PR. Differences must be explainable by flat hit + normals / card-res albedo (Tier 2 closes those). +- Toggling F9 off restores the raster path with no residue. + +## Status (2026-07-13, AMD 760M / DX12+DXC) + +- **M1 passed**: `BLOOM_PT_DEBUG=5` magenta over the full frame, HUD + + translucent water compositing on top, no validation errors. +- Depth convention confirmed standard-Z (`BLOOM_PT_DEBUG=1`: near dark, + far bright) — `is_sky(depth >= 0.9999999)` is correct as written. +- **M2 passed**: full render at brightness parity with raster; static + title camera at ~17 s vs ~50 s shows clearly reduced grass/canopy + noise (convergence, not mere difference). Fresh-accumulation noise + visible as expected right after a camera cut. +- Live mode switching verified: F9 → RT (EMA) mid-game, F9 → off + restores the raster path bit-clean (SSGI/SSR resume, 13 → 46 fps). +- **M3 partially open**: point-light NEE + building-base colour bleed + + `bloom-reference` RMSE not yet measured (needs a lit world + fixed + exposure; folded into PT-2/PT-5 verification). + +Implementation notes discovered on the way: + +- `current_vp_matrix` carries the TAA Halton jitter (~1e-3 in the proj + Z-coupling slots). The accumulation-reset check MUST compare an + unjittered VP (`current_proj_matrix_unjittered × view`) or mode 1 + pins at 1 sample forever. The *jittered* `inv_vp` still feeds the + kernel — primary rays must match the jittered G-buffer depth, and + accumulating across jitters is free AA. +- TAA runs downstream of PT, so converged PT output gets a second + temporal filter (soft distant foliage). Acceptable for Tier 1; + PT-3 owns the proper answer (bypass/replace TAA under PT). +- Ray queries treat leaf-card cutouts as solid quads on bounce rays + (slight over-darkening under canopies). Primary visibility comes + from the G-buffer so silhouettes are unaffected. PT-2 material at + hits can alpha-test via `RAY_FLAG_NO_OPAQUE` + candidate loop if it + proves visible. +- Batch verification: `BLOOM_PT`/`BLOOM_PT_DEBUG` env seed the renderer + directly. Careful with F12-capture scripts: the title screen treats + the posted F12 as "press any key" and starts the game. + +## Known honest limits (by design, closed by PT-2) + +Flat per-mesh hit normals; card-resolution albedo at hits; no specular +bounce; no glass; sky at miss is the analytic Lumen sky, not the LUT sky. diff --git a/native/shared/src/renderer/formats.rs b/native/shared/src/renderer/formats.rs index 076050b..78febf3 100644 --- a/native/shared/src/renderer/formats.rs +++ b/native/shared/src/renderer/formats.rs @@ -92,9 +92,13 @@ pub(super) fn create_hdr_rt(device: &wgpu::Device, width: u32, height: u32) -> ( // Phase 4b adds COPY_SRC so the translucent-pass scheduler // can snapshot hdr_rt → a SceneColor transient via // copy_texture_to_texture before refractive draws run. + // + // STORAGE_BINDING: the path-trace megakernel (PT-1) writes the + // traced scene colour directly into hdr_rt from compute. usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::TEXTURE_BINDING - | wgpu::TextureUsages::COPY_SRC, + | wgpu::TextureUsages::COPY_SRC + | wgpu::TextureUsages::STORAGE_BINDING, view_formats: &[], }); let view = texture.create_view(&wgpu::TextureViewDescriptor::default()); diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 4488c08..a1bb1dd 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -8,6 +8,7 @@ mod hiz; mod occlusion; mod ssr_pass; mod ssgi_pass; +mod pt_pass; mod shadow_pass; mod model_draw; mod planar_pass; @@ -358,6 +359,29 @@ struct ProbeTraceParams { wsrc_cascades: [[f32; 4]; 3], } +/// Uniform for the path-tracing megakernel — WGSL mirror `PtParams` in +/// shaders/pt.rs. Field-for-field, 16-byte aligned, 672 bytes. +#[repr(C)] +#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] +struct PtParamsCpu { + inv_vp: [[f32; 4]; 4], + /// xyz camera world pos, w unused. + cam_pos: [f32; 4], + /// xyz unit vector toward the sun, w unused. + sun_dir: [f32; 4], + /// rgb premultiplied by intensity, w unused. + sun_color: [f32; 4], + /// rgb ambient-derived sky tint, w unused. + sky_color: [f32; 4], + /// x=width, y=height, z=frame_index, w=accum_count. + size: [u32; 4], + /// x=mode (1 progressive / 2 realtime), y=max_bounces, + /// z=point_light_count (≤16), w=debug view (BLOOM_PT_DEBUG). + cfg: [f32; 4], + /// 16 lights × (pos_range vec4, color_int vec4). + lights: [[f32; 4]; 32], +} + #[repr(C)] #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] struct ProbeTemporalParams { @@ -1215,6 +1239,28 @@ pub struct Renderer { /// SW cache shape. probe_trace_hw_bg_cache: [Option; 2], + // --- Path tracing (PT-1, docs/pt/pt-roadmap.md) --- + /// Megakernel pipeline; None when the adapter lacks ray query + /// (same gate as the HW probe trace — there is no SW fallback). + pub pt_pipeline: Option, + pub pt_layout: Option, + /// Rebuilt lazily; nulled on resize and on instance-data rebuild. + pt_bg: Option, + pt_uniform_buffer: wgpu::Buffer, + /// rgba32float-equivalent accumulation (vec4 per pixel) as a storage + /// buffer — rgba32float storage *textures* lack read_write in core + /// WGSL. Lazily (re)created at render extent. + pt_accum_buffer: Option, + /// Samples accumulated at the current view; 0 = history invalid. + pt_accum_count: u32, + /// VP matrix of the last accumulated frame; movement beyond epsilon + /// resets accumulation (mode 1) — mode 2 relies on EMA instead. + pt_prev_vp: [[f32; 4]; 4], + /// TLAS version last traced; a rebuild invalidates accumulation. + pt_last_tlas_version: u64, + /// BLOOM_PT_DEBUG view selector, forwarded to the kernel via cfg.w. + pt_debug: f32, + /// Ticket 014 V3 — SW SDF sphere-trace pipeline + layout. /// Active whenever the scene clipmap has been baked; chosen at /// dispatch time over the Hi-Z SW fallback when available. @@ -4915,6 +4961,123 @@ impl Renderer { (None, None) }; + // --- Path-tracing megakernel (PT-1, docs/pt/pt-roadmap.md) --- + // Same gate as the HW probe trace: no ray query, no path tracer — + // and deliberately no software fallback (a CPU-speed path trace is + // seconds per frame, worse than useless in any shipped mode). + let (pt_pipeline, pt_layout) = if hw_rt_enabled { + let pt_source = format!("enable wgpu_ray_query;\n{}", PT_KERNEL_WGSL); + let pt_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("pt_kernel_shader"), + source: wgpu::ShaderSource::Wgsl(pt_source.into()), + }); + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("pt_layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::AccelerationStructure { + vertex_return: false, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 3, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 4, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 5, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 6, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 7, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 8, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 9, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: HDR_FORMAT, + view_dimension: wgpu::TextureViewDimension::D2, + }, count: None, + }, + ], + }); + let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("pt_pl_layout"), + bind_group_layouts: &[Some(&layout)], + immediate_size: 0, + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("pt_pipeline"), + layout: Some(&pl), + module: &pt_shader, entry_point: Some("cs_main"), + compilation_options: Default::default(), cache: None, + }); + (Some(pipeline), Some(layout)) + } else { + (None, None) + }; + let pt_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_uniform"), + size: std::mem::size_of::() as u64, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + // Debug view selector for the kernel (see shaders/pt.rs header). + // Env read happens once at device init, which is fine: it is a + // developer-only diagnostic, not a runtime setting. + let pt_debug = std::env::var("BLOOM_PT_DEBUG") + .ok() + .and_then(|v| v.parse::().ok()) + .unwrap_or(0.0); + // --- Ticket 014 V3: SW SDF sphere-trace pipeline --- // Always built. At dispatch time we pick SDF over Hi-Z when // `scene_sdf_clipmap_built` is true; HW (when available) @@ -6477,7 +6640,14 @@ impl Renderer { ssgi_intensity: 1.0, ssgi_radius: 20.0, ssgi_enabled: true, - pt_mode: 0, + // BLOOM_PT env seeds the initial mode for headless/batch + // verification runs (keypresses don't reach batch runs); + // the game overrides it live via set_path_tracing. + pt_mode: std::env::var("BLOOM_PT") + .ok() + .and_then(|v| v.parse::().ok()) + .unwrap_or(0) + .min(2), ssgi_backend_logged: None, gi_scene_avg_albedo: [0.35, 0.35, 0.35], probe_grid_w, @@ -6504,6 +6674,15 @@ impl Renderer { probe_trace_hw_pipeline, probe_trace_hw_layout, probe_trace_hw_bg_cache: [None, None], + pt_pipeline, + pt_layout, + pt_bg: None, + pt_uniform_buffer, + pt_accum_buffer: None, + pt_accum_count: 0, + pt_prev_vp: [[0.0; 4]; 4], + pt_last_tlas_version: 0, + pt_debug, probe_trace_sdf_pipeline, probe_trace_sdf_layout, probe_trace_sdf_bg_cache: [None, None], @@ -6999,6 +7178,12 @@ impl Renderer { self.probe_trace_sdf_bg_cache = [None, None]; self.probe_temporal_bg_cache = [None, None]; self.probe_resolve_bg_cache = [None, None]; + // PT-1 — the BG references hdr_rt/depth views and the + // accumulation buffer is per-pixel-sized; both die with + // the old extent. + self.pt_bg = None; + self.pt_accum_buffer = None; + self.pt_accum_count = 0; self.hiz_linearize_bg_cache = None; self.occlusion.invalidate_bindings(); for slot in self.hiz_downsample_bg_cache.iter_mut() { @@ -7684,6 +7869,8 @@ impl Renderer { // V14 — WSRC HW bake bg also references the TLAS + // instance_data buffer; invalidate on resize. self.wsrc_bake_hw_bg_cache = None; + // PT-1 — same buffer bound at group 0 binding 2. + self.pt_bg = None; resized = true; } @@ -7787,6 +7974,8 @@ impl Renderer { self.probe_trace_hw_bg_cache = [None, None]; // V14 — same reason as the resize path above. self.wsrc_bake_hw_bg_cache = None; + // PT-1 — TLAS bound at group 0 binding 1. + self.pt_bg = None; } // Populate TLAS instance slots. Clear stale entries from prior @@ -9600,6 +9789,18 @@ impl Renderer { ]) .with_after(&["shadow", "froxel_assign"]), ); + g.push( + PassNode::new("pt", Box::new(|c: &mut FrameCtx2| { + // No-op unless pt_active() (the method gates). Sits + // between the opaque scene and translucency so water / + // glass still composite over the traced opaques, and + // sky pixels (never written by the kernel) keep the + // raster sky. + c.r.record_pt_pass(c.encoder, c.profiler, c.surf.0, c.surf.1); + })) + .with_reads(&[PassInput::SceneDepth]) + .with_after(&["hdr_scene"]), + ); g.push( PassNode::new("translucent", Box::new(|c: &mut FrameCtx2| { c.r.record_translucent_pass(c.encoder, c.profiler); @@ -9607,7 +9808,7 @@ impl Renderer { // Reads the opaque HDR + depth and alpha-blends back into // HdrColor; the pin (not a second HdrColor write) keeps a // single declared writer per resource. - .with_after(&["hdr_scene"]), + .with_after(&["pt"]), ); g.push( PassNode::new("hiz_build", Box::new(|c: &mut FrameCtx2| { diff --git a/native/shared/src/renderer/postfx_chain.rs b/native/shared/src/renderer/postfx_chain.rs index 2f46c38..8a4845b 100644 --- a/native/shared/src/renderer/postfx_chain.rs +++ b/native/shared/src/renderer/postfx_chain.rs @@ -182,7 +182,10 @@ impl Renderer { encoder: &mut wgpu::CommandEncoder, ) { // Composite input views (were locals in end_frame_with_scene). - let ssr_composite_view = if self.ssr_enabled { + // PT-1: while path tracing, the SSR passes are skipped entirely, + // so history is stale — route compose to ssr_rt, which the march + // else-branch keeps cleared to transparent black. + let ssr_composite_view = if self.ssr_enabled && !self.pt_active() { &self.ssr_history_views[self.ssr_history_idx] } else { &self.ssr_rt_view diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs new file mode 100644 index 0000000..4b590f6 --- /dev/null +++ b/native/shared/src/renderer/pt_pass.rs @@ -0,0 +1,175 @@ +//! PT-1 — progressive path-trace megakernel dispatch +//! (docs/pt/PT-1-progressive-megakernel.md). Replaces the lit opaque +//! scene colour in hdr_rt when path tracing is active; sky pixels are +//! left untouched so the raster sky/clouds survive, and translucency +//! still composites on top afterwards. + +use super::*; + +impl Renderer { + pub(super) fn record_pt_pass( + &mut self, + encoder: &mut wgpu::CommandEncoder, + profiler: &mut crate::profiler::Profiler, + surf_w: u32, + surf_h: u32, + ) { + if !self.pt_active() { + // Leaving PT (or never entering it) invalidates history so + // re-enabling starts a fresh accumulation, not a stale one. + self.pt_accum_count = 0; + return; + } + // Same readiness gate as the HW probe trace: first frames before + // any geometry is committed have no TLAS / instance data yet. + if self.pt_pipeline.is_none() + || self.tlas.is_none() + || self.tlas_instance_data_buffer.is_none() + { + self.pt_accum_count = 0; + return; + } + + // ---- accumulation validity ---- + // Any camera motion beyond epsilon restarts progressive + // accumulation (mode 1). Mode 2 (realtime) ignores the reset — + // its EMA is designed to absorb motion — but tracking prev_vp + // costs nothing and keeps one code path. + // + // Compared UNJITTERED: current_vp_matrix carries the TAA Halton + // nudge (~1e-3 in the proj Z-coupling slots), which would read + // as motion every frame and pin the accumulator at 1 sample. + // The jittered inv_vp still goes to the kernel — primary rays + // must match the jittered G-buffer depth, and accumulating + // across jitters is free anti-aliasing. + let vp_unjittered = mat4_multiply( + self.current_proj_matrix_unjittered, + self.current_view_matrix, + ); + let mut moved = false; + for r in 0..4 { + for c in 0..4 { + if (vp_unjittered[r][c] - self.pt_prev_vp[r][c]).abs() > 1e-5 { + moved = true; + } + } + } + if moved && self.pt_mode == 1 { + self.pt_accum_count = 0; + } + self.pt_prev_vp = vp_unjittered; + // Geometry changed under the accumulated image (door opened, + // enemy died) → history is a lie, restart. + if self.tlas_built_version != self.pt_last_tlas_version { + self.pt_last_tlas_version = self.tlas_built_version; + self.pt_accum_count = 0; + } + + // ---- accumulation buffer (vec4 per pixel) ---- + let needed = (surf_w as u64) * (surf_h as u64) * 16; + let recreate = match &self.pt_accum_buffer { + Some(b) => b.size() != needed, + None => true, + }; + if recreate { + self.pt_accum_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_accum"), + size: needed, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + self.pt_bg = None; + self.pt_accum_count = 0; + } + + // ---- uniforms ---- + // Sun / sky derivation matches record_ssgi_passes exactly so PT + // brightness lines up with the raster + GI frame it replaces. + let ld = self.lighting_uniforms.light_dir; + let sun_inv_len = 1.0 / (ld[0] * ld[0] + ld[1] * ld[1] + ld[2] * ld[2]).sqrt().max(1e-4); + let sun_intensity = ld[3].max(0.0); + let lc = self.lighting_uniforms.light_color; + let amb = self.lighting_uniforms.ambient; + let sky_intensity = amb[3].max(0.0); + + let light_count = (self.lighting_uniforms.point_light_count[0] as usize).min(16); + let mut lights = [[0.0f32; 4]; 32]; + for i in 0..light_count { + let pl = &self.lighting_uniforms.point_lights[i]; + lights[i * 2] = pl.position; // xyz + range + lights[i * 2 + 1] = pl.color; // rgb + intensity + } + + let max_bounces = if self.pt_mode == 2 { 2.0 } else { 8.0 }; + let cam = self.current_camera_pos; + let params = PtParamsCpu { + inv_vp: self.current_inv_vp_matrix, + cam_pos: [cam[0], cam[1], cam[2], 0.0], + sun_dir: [ + -ld[0] * sun_inv_len, + -ld[1] * sun_inv_len, + -ld[2] * sun_inv_len, + 0.0, + ], + sun_color: [ + lc[0] * sun_intensity, + lc[1] * sun_intensity, + lc[2] * sun_intensity, + 0.0, + ], + sky_color: [ + amb[0] * sky_intensity, + amb[1] * sky_intensity, + amb[2] * sky_intensity, + 0.0, + ], + size: [surf_w, surf_h, self.taa_frame_index, self.pt_accum_count], + cfg: [ + self.pt_mode as f32, + max_bounces, + light_count as f32, + self.pt_debug, + ], + lights, + }; + self.queue.write_buffer(&self.pt_uniform_buffer, 0, bytemuck::bytes_of(¶ms)); + + // ---- bind group (lazy; nulled on resize / TLAS or instance + // buffer recreation) ---- + if self.pt_bg.is_none() { + let tlas = self.tlas.as_ref().unwrap(); + self.pt_bg = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_bg"), + layout: self.pt_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_uniform_buffer.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: tlas.as_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: self.tlas_instance_data_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.material_rt_view) }, + // Raw albedo atlas, NOT the pre-lit radiance atlas the + // GI probe trace uses — PT computes its own lighting at + // hits; radiance would double-count. + wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.mesh_card_atlas_view) }, + wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::Sampler(&self.mesh_card_atlas_sampler) }, + wgpu::BindGroupEntry { binding: 8, resource: self.pt_accum_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + ], + })); + } + + // ---- dispatch ---- + { + let ts = profiler.compute_pass_timestamp_writes("pt_pass"); + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("pt_pass"), + timestamp_writes: ts, + }); + pass.set_pipeline(self.pt_pipeline.as_ref().unwrap()); + pass.set_bind_group(0, self.pt_bg.as_ref().unwrap(), &[]); + pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); + } + self.pt_accum_count = self.pt_accum_count.saturating_add(1); + } +} diff --git a/native/shared/src/renderer/shaders/mod.rs b/native/shared/src/renderer/shaders/mod.rs index 6092ed3..312e766 100644 --- a/native/shared/src/renderer/shaders/mod.rs +++ b/native/shared/src/renderer/shaders/mod.rs @@ -13,5 +13,7 @@ mod gi; pub(super) use gi::{CARD_CAPTURE_WGSL, SDF_BAKE_WGSL, SDF_CLIPMAP_BAKE_WGSL, CARD_LIGHT_WGSL, WSRC_BAKE_WGSL, WSRC_BAKE_HW_WGSL}; mod ssgi; pub(super) use ssgi::{PROBE_HELPERS_WGSL, SSGI_PROBE_PLACE_WGSL, SSGI_PROBE_TRACE_SW_WGSL, SSGI_PROBE_TRACE_HW_WGSL, SSGI_PROBE_TRACE_SDF_WGSL, SSGI_PROBE_TEMPORAL_WGSL, SSGI_PROBE_RESOLVE_WGSL, SSR_TEMPORAL_SHADER_WGSL, SSR_SHADER_WGSL}; +mod pt; +pub(super) use pt::PT_KERNEL_WGSL; mod post; pub(super) use post::{BLOOM_SHADER_WGSL, DOF_SHADER_WGSL, MOTION_BLUR_SHADER_WGSL, SSS_SHADER_WGSL, SCENE_COMPOSE_SHADER_WGSL, TAA_SHADER_WGSL, EXPOSURE_SHADER_WGSL, COMPOSITE_SHADER_WGSL, UPSCALE_SHADER_WGSL, RCAS_SHADER_WGSL}; diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs new file mode 100644 index 0000000..ad1b977 --- /dev/null +++ b/native/shared/src/renderer/shaders/pt.rs @@ -0,0 +1,421 @@ +//! Path-tracing megakernel (docs/pt/pt-roadmap.md, ticket PT-1). +//! +//! One compute kernel, one ray budget per pixel per frame. Primary hits come +//! from the G-buffer (depth + albedo + material MRTs — free, and sharper than +//! traced primaries); bounce and shadow rays go through the same TLAS the +//! Lumen HW probe trace uses. Hit shading at bounces reads the mesh-card +//! ALBEDO atlas (not the pre-lit radiance atlas: a path tracer computes its +//! own lighting at every vertex of the path — sampling pre-lit cards would +//! bake Lumen's direct light into ours twice). +//! +//! Radiometric convention: light intensities are treated as π-premultiplied, +//! i.e. diffuse contribution is `albedo * L * NdotL` with no 1/π — matching +//! the raster shader (core.rs point-light loop has no 1/π either), so +//! toggling PT on/off does not jump scene brightness. bloom-reference +//! comparisons account for this in scene config (see the PT-1 ticket). +//! +//! Sky pixels are never written: the raster sky/cloud passes already drew +//! them, and PT replacing a procedural cloud deck with an analytic gradient +//! would be a downgrade. PT owns geometry pixels only. The translucent pass +//! runs AFTER this kernel, so water and glass composite over path-traced +//! opaques exactly as they do over raster ones. +//! +//! Debug modes (uniform cfg.w, set via BLOOM_PT_DEBUG): +//! 1 = raw depth visualised 2 = reconstructed world normals +//! 3 = G-buffer albedo 4 = sun shadow-ray visibility +//! 5 = solid magenta (pipeline probe — proves dispatch + write path) + +pub(in crate::renderer) const PT_KERNEL_WGSL: &str = r#" +struct PtLight { + pos_range: vec4, // xyz world position, w = range + color_int: vec4, // rgb color, w = intensity +}; + +struct PtParams { + inv_vp: mat4x4, + cam_pos: vec4, // xyz camera world pos + sun_dir: vec4, // xyz unit vector toward the sun + sun_color: vec4, // rgb premultiplied by intensity + sky_color: vec4, // rgb ambient-derived sky tint + size: vec4, // x=width, y=height, z=frame_index, w=accum_count + cfg: vec4, // x=mode(1|2), y=max_bounces, z=point_light_count, w=debug + lights: array, +}; + +// Layout mirror of the Lumen instance data (ssgi.rs) — same buffer. +struct InstanceGiData { + albedo: vec3, + emissive_luma: f32, + normal_ws: vec3, + _pad0: f32, + card_slot: vec4, + card_aabb_min: vec4, + card_aabb_max: vec4, + world_aabb_min: vec4, + world_aabb_max: vec4, +}; + +@group(0) @binding(0) var u: PtParams; +@group(0) @binding(1) var accel: acceleration_structure; +@group(0) @binding(2) var instance_data: array; +@group(0) @binding(3) var depth_tex: texture_depth_2d; +@group(0) @binding(4) var albedo_tex: texture_2d; +@group(0) @binding(5) var material_tex: texture_2d; +@group(0) @binding(6) var card_albedo_atlas: texture_2d; +@group(0) @binding(7) var card_samp: sampler; +@group(0) @binding(8) var accum: array>; +@group(0) @binding(9) var out_hdr: texture_storage_2d; + +// ---- RNG: PCG, one stream per (pixel, frame) -------------------------------- + +var rng_state: u32; + +fn rng_seed(px: vec2, frame: u32) { + var h = px.x * 374761393u + px.y * 668265263u + frame * 2654435761u; + h = (h ^ (h >> 13u)) * 1274126177u; + rng_state = h ^ (h >> 16u); +} + +fn rand_f() -> f32 { + // PCG-XSH-RR step. + let old = rng_state; + rng_state = old * 747796405u + 2891336453u; + let word = ((old >> ((old >> 28u) + 4u)) ^ old) * 277803737u; + let out = (word >> 22u) ^ word; + return f32(out) * 2.3283064e-10; // / 2^32 +} + +fn rand_2f() -> vec2 { return vec2(rand_f(), rand_f()); } + +// ---- Geometry reconstruction ------------------------------------------------- + +fn world_at(px: vec2, depth: f32) -> vec3 { + let dims = vec2(f32(u.size.x), f32(u.size.y)); + let uv = (vec2(px) + vec2(0.5)) / dims; + let ndc = vec4(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, depth, 1.0); + let w = u.inv_vp * ndc; + return w.xyz / w.w; +} + +fn depth_at(px: vec2) -> f32 { + let clamped = clamp(px, vec2(0), vec2(i32(u.size.x) - 1, i32(u.size.y) - 1)); + return textureLoad(depth_tex, clamped, 0); +} + +fn is_sky(depth: f32) -> bool { + // Depth-buffer far plane. If the projection turns out reversed-Z the + // BLOOM_PT_DEBUG=1 depth view makes it obvious in one screenshot; flip + // here if geometry reads bright and sky reads dark. + return depth >= 0.9999999; +} + +// Screen-space normal from depth: reconstruct neighbours, take the tighter +// derivative on each axis so depth discontinuities don't smear normals +// across silhouettes. +fn normal_from_depth(px: vec2, p_center: vec3) -> vec3 { + let d_l = depth_at(px + vec2(-1, 0)); + let d_r = depth_at(px + vec2(1, 0)); + let d_u = depth_at(px + vec2(0, -1)); + let d_d = depth_at(px + vec2(0, 1)); + let d_c = depth_at(px); + + var ddx: vec3; + if (abs(d_l - d_c) < abs(d_r - d_c)) { + ddx = p_center - world_at(px + vec2(-1, 0), d_l); + } else { + ddx = world_at(px + vec2(1, 0), d_r) - p_center; + } + var ddy: vec3; + if (abs(d_u - d_c) < abs(d_d - d_c)) { + ddy = p_center - world_at(px + vec2(0, -1), d_u); + } else { + ddy = world_at(px + vec2(0, 1), d_d) - p_center; + } + var n = cross(ddy, ddx); + let len = length(n); + if (len < 1e-8) { return vec3(0.0, 1.0, 0.0); } + n = n / len; + // Face the camera: a G-buffer surface always does. + if (dot(n, u.cam_pos.xyz - p_center) < 0.0) { n = -n; } + return n; +} + +// ---- Sampling helpers ---------------------------------------------------------- + +// Branchless ONB (Duff et al. 2017). +fn onb(n: vec3) -> mat3x3 { + let s = select(-1.0, 1.0, n.z >= 0.0); + let a = -1.0 / (s + n.z); + let b = n.x * n.y * a; + let t = vec3(1.0 + s * n.x * n.x * a, s * b, -s * n.x); + let bt = vec3(b, s + n.y * n.y * a, -n.y); + return mat3x3(t, bt, n); +} + +fn cosine_sample(n: vec3, r: vec2) -> vec3 { + let phi = 6.2831853 * r.x; + let sr = sqrt(r.y); + let local = vec3(cos(phi) * sr, sin(phi) * sr, sqrt(max(0.0, 1.0 - r.y))); + return normalize(onb(n) * local); +} + +// Uniform direction in the solar cone (half-angle 0.265 deg -> soft shadows). +fn sun_cone_sample(r: vec2) -> vec3 { + let cos_max = 0.9999893; + let cos_t = mix(cos_max, 1.0, r.x); + let sin_t = sqrt(max(0.0, 1.0 - cos_t * cos_t)); + let phi = 6.2831853 * r.y; + let local = vec3(cos(phi) * sin_t, sin(phi) * sin_t, cos_t); + return normalize(onb(u.sun_dir.xyz) * local); +} + +// Sky radiance for a miss. Analytic horizon-to-zenith gradient off the same +// ambient-derived tint Lumen's traces use; the sun disc is deliberately +// absent (the sun is sampled by NEE only, so it cannot be counted twice). +fn sky_radiance(dir: vec3) -> vec3 { + let t = clamp(dir.y * 0.5 + 0.5, 0.0, 1.0); + return u.sky_color.rgb * mix(0.45, 1.35, t); +} + +// ---- Ray casts ------------------------------------------------------------------ + +fn occluded(origin: vec3, dir: vec3, max_t: f32) -> bool { + var rq: ray_query; + rayQueryInitialize(&rq, accel, RayDesc(0u, 0xFFu, 0.001, max_t, origin, dir)); + loop { + if (!rayQueryProceed(&rq)) { break; } + } + let hit = rayQueryGetCommittedIntersection(&rq); + return hit.kind != RAY_QUERY_INTERSECTION_NONE; +} + +// ---- Hit shading: card albedo ----------------------------------------------------- + +// Same signed-axis card projection as the Lumen HW trace (ssgi.rs), but +// sampling the raw ALBEDO atlas. Falls back to the flat instance albedo when +// the mesh has no captured card. +fn albedo_at_hit( + inst: InstanceGiData, + hit_os: vec3, + dir_ws: vec3, +) -> vec3 { + if (inst.card_slot.w <= 0.5) { + return inst.albedo; + } + let abs_d = abs(dir_ws); + var axis_idx: u32 = 0u; + if (abs_d.y >= abs_d.x && abs_d.y >= abs_d.z) { + axis_idx = 2u; + } else if (abs_d.z >= abs_d.x) { + axis_idx = 4u; + } + var signed_axis: u32 = axis_idx; + if (axis_idx == 0u && dir_ws.x > 0.0) { signed_axis = 1u; } + else if (axis_idx == 2u && dir_ws.y > 0.0) { signed_axis = 3u; } + else if (axis_idx == 4u && dir_ws.z > 0.0) { signed_axis = 5u; } + + let first_slot = u32(inst.card_slot.x); + let slot = first_slot + signed_axis; + let slot_x = slot % 64u; + let slot_y = slot / 64u; + + let bmin = inst.card_aabb_min.xyz; + let bmax = inst.card_aabb_max.xyz; + var u_os: f32; var v_os: f32; + var u_lo: f32; var u_hi: f32; var v_lo: f32; var v_hi: f32; + var u_flip: f32 = 1.0; + if (signed_axis == 0u || signed_axis == 1u) { + u_os = hit_os.y; v_os = hit_os.z; + u_lo = bmin.y; u_hi = bmax.y; v_lo = bmin.z; v_hi = bmax.z; + if (signed_axis == 1u) { u_flip = -1.0; } + } else if (signed_axis == 2u || signed_axis == 3u) { + u_os = hit_os.x; v_os = hit_os.z; + u_lo = bmin.x; u_hi = bmax.x; v_lo = bmin.z; v_hi = bmax.z; + if (signed_axis == 3u) { u_flip = -1.0; } + } else { + u_os = hit_os.x; v_os = hit_os.y; + u_lo = bmin.x; u_hi = bmax.x; v_lo = bmin.y; v_hi = bmax.y; + if (signed_axis == 5u) { u_flip = -1.0; } + } + var u_norm = clamp((u_os - u_lo) / max(u_hi - u_lo, 1e-4), 0.0, 1.0); + let v_norm = clamp((v_os - v_lo) / max(v_hi - v_lo, 1e-4), 0.0, 1.0); + if (u_flip < 0.0) { u_norm = 1.0 - u_norm; } + + let slot_size_uv = 1.0 / 64.0; + let texel_in_slot = slot_size_uv / 64.0; + let slot_u0 = f32(slot_x) * slot_size_uv + texel_in_slot; + let slot_v0 = f32(slot_y) * slot_size_uv + texel_in_slot; + let slot_span = slot_size_uv - 2.0 * texel_in_slot; + let atlas_uv = vec2(slot_u0 + u_norm * slot_span, slot_v0 + v_norm * slot_span); + return textureSampleLevel(card_albedo_atlas, card_samp, atlas_uv, 0.0).rgb; +} + +// ---- Next-event estimation --------------------------------------------------------- + +// Direct light at a surface point: sun through the solar cone + one point +// light chosen uniformly (contribution / pdf). Game-radiometry convention: +// no 1/pi (see file header). +fn direct_light(p: vec3, n: vec3, alb: vec3) -> vec3 { + var lit = vec3(0.0); + + let sd = sun_cone_sample(rand_2f()); + let ndl = dot(n, sd); + if (ndl > 0.0 && !occluded(p, sd, 1000.0)) { + lit += u.sun_color.rgb * ndl; + } + + let count = u32(u.cfg.z); + if (count > 0u) { + let pick = min(u32(rand_f() * f32(count)), count - 1u); + let l = u.lights[pick]; + let to_l = l.pos_range.xyz - p; + let d = length(to_l); + let range = l.pos_range.w; + if (d < range && d > 1e-3) { + let dir = to_l / d; + let ndl2 = dot(n, dir); + if (ndl2 > 0.0 && !occluded(p, dir, d - 0.02)) { + // Raster-parity falloff: (1 - d/range)^2, core.rs. + let att = 1.0 - d / range; + lit += l.color_int.rgb * l.color_int.w * ndl2 * att * att * f32(count); + } + } + } + return alb * lit; +} + +// ---- Main ----------------------------------------------------------------------------- + +@compute @workgroup_size(8, 8, 1) +fn cs_main(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= u.size.x || gid.y >= u.size.y) { return; } + let px = vec2(i32(gid.x), i32(gid.y)); + rng_seed(gid.xy, u.size.z); + + let debug = u.cfg.w; + if (debug == 5.0) { + textureStore(out_hdr, px, vec4(1.0, 0.0, 1.0, 1.0)); + return; + } + + let depth = depth_at(px); + if (debug == 1.0) { + textureStore(out_hdr, px, vec4(vec3(depth), 1.0)); + return; + } + + if (is_sky(depth)) { + // Leave the raster sky/clouds untouched. + return; + } + + let p0 = world_at(px, depth); + let n0 = normal_from_depth(px, p0); + let albedo0 = textureLoad(albedo_tex, px, 0).rgb; + + if (debug == 2.0) { + textureStore(out_hdr, px, vec4(n0 * 0.5 + 0.5, 1.0)); + return; + } + if (debug == 3.0) { + textureStore(out_hdr, px, vec4(albedo0, 1.0)); + return; + } + if (debug == 4.0) { + let sd = sun_cone_sample(rand_2f()); + let vis = select(0.0, 1.0, !occluded(p0 + n0 * 0.02, sd, 1000.0)); + textureStore(out_hdr, px, vec4(vec3(vis), 1.0)); + return; + } + + // ---- one path sample -------------------------------------------------- + + var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0); + var throughput = albedo0; + var origin = p0 + n0 * 0.02; + var n_cur = n0; + + let max_bounces = u32(u.cfg.y); + for (var b = 0u; b < max_bounces; b = b + 1u) { + let dir = cosine_sample(n_cur, rand_2f()); + + var rq: ray_query; + rayQueryInitialize(&rq, accel, RayDesc(0u, 0xFFu, 0.001, 500.0, origin, dir)); + loop { + if (!rayQueryProceed(&rq)) { break; } + } + let hit = rayQueryGetCommittedIntersection(&rq); + + if (hit.kind == RAY_QUERY_INTERSECTION_NONE) { + radiance += throughput * sky_radiance(dir); + break; + } + + let inst = instance_data[hit.instance_custom_data]; + let hit_ws = origin + dir * hit.t; + let hit_os = (hit.world_to_object * vec4(hit_ws, 1.0)).xyz; + let alb_hit = albedo_at_hit(inst, hit_os, dir); + + // Hit normal: the flat per-instance normal (Tier-1 honesty note in + // the roadmap; PT-2 replaces this with interpolated vertex normals). + var n_hit = inst.normal_ws; + let n_len = length(n_hit); + if (n_len < 1e-4) { + n_hit = -dir; + } else { + n_hit = n_hit / n_len; + if (dot(n_hit, dir) > 0.0) { n_hit = -n_hit; } + } + + // Emissive surfaces radiate; matches the Lumen fallback semantics + // (albedo * emissive_luma). + radiance += throughput * inst.albedo * inst.emissive_luma; + + let hit_p = hit_ws + n_hit * 0.02; + radiance += throughput * direct_light(hit_p, n_hit, alb_hit); + + throughput *= alb_hit; + origin = hit_p; + n_cur = n_hit; + + // Russian roulette from the third bounce. + if (b >= 2u) { + let q = clamp(max(throughput.r, max(throughput.g, throughput.b)), 0.05, 0.95); + if (rand_f() > q) { break; } + throughput /= q; + } + } + + // NaN/Inf guard + firefly cap so one bad sample cannot poison the + // accumulator forever. + if (radiance.r != radiance.r || radiance.g != radiance.g || radiance.b != radiance.b) { + radiance = vec3(0.0); + } + let luma = dot(radiance, vec3(0.2126, 0.7152, 0.0722)); + if (luma > 32.0) { radiance *= 32.0 / luma; } + + // ---- accumulate --------------------------------------------------------- + + let idx = gid.y * u.size.x + gid.x; + let mode = u.cfg.x; + var prev = accum[idx]; + if (u.size.w == 0u) { prev = vec4(0.0); } + + var out: vec3; + if (mode >= 2.0) { + // Realtime: rolling exponential window. PT-3 replaces this with + // motion-reprojected accumulation + a-trous. + let alpha = select(0.125, 1.0, u.size.w == 0u); + out = mix(prev.rgb, radiance, alpha); + accum[idx] = vec4(out, 1.0); + } else { + // Progressive: plain running sum; count lives on the CPU. + let sum = prev.rgb + radiance; + let n = f32(u.size.w) + 1.0; + accum[idx] = vec4(sum, n); + out = sum / n; + } + + textureStore(out_hdr, px, vec4(out, 1.0)); +} +"#; diff --git a/native/shared/src/renderer/ssgi_pass.rs b/native/shared/src/renderer/ssgi_pass.rs index 2e18295..f4f5e54 100644 --- a/native/shared/src/renderer/ssgi_pass.rs +++ b/native/shared/src/renderer/ssgi_pass.rs @@ -27,7 +27,11 @@ impl Renderer { let write_idx = self.probe_history_idx; let prev_idx = 1 - write_idx; - if self.ssgi_enabled { + // PT-1: while the path tracer owns the frame its output already + // contains full GI — probe SSGI would burn ~2ms to be composited + // over by nothing (the else-branch clear keeps compose additive- + // safe either way). + if self.ssgi_enabled && !self.pt_active() { let p00 = self.current_proj_matrix[0][0]; let p11 = self.current_proj_matrix[1][1]; let p20 = self.current_proj_matrix[2][0]; diff --git a/native/shared/src/renderer/ssr_pass.rs b/native/shared/src/renderer/ssr_pass.rs index 568cdd1..a56f992 100644 --- a/native/shared/src/renderer/ssr_pass.rs +++ b/native/shared/src/renderer/ssr_pass.rs @@ -15,7 +15,9 @@ impl Renderer { // ============================================================ // SSR: view-space ray march of the depth buffer + HDR sample. // ============================================================ - if self.ssr_enabled { + // PT-1: skipped while the path tracer owns the frame — it marches + // the raster-lit HDR, which PT has already overwritten. + if self.ssr_enabled && !self.pt_active() { let inv_proj = self.current_inv_proj_matrix; // EN-021 — view→world rotation for the env-miss fallback: the // transpose of the view matrix's 3×3 (rigid view ⇒ inverse @@ -136,7 +138,8 @@ impl Renderer { // history. Compose then reads ssr_history[cur] instead of // ssr_rt. // ============================================================ - if self.ssr_enabled { + // PT-1: same gate as the march — no fresh rays, nothing to blend. + if self.ssr_enabled && !self.pt_active() { let prev_idx = 1 - self.ssr_history_idx; let cur_idx = self.ssr_history_idx; From 74820ee0f1d902b2f8e57c77760093001c559e46 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 18:24:49 +0200 Subject: [PATCH 03/16] feat(pt): PT-2 real hit shading + fix transposed inv_vp ray generation Geometry megabuffer (concatenated Vertex3D + index words, per-instance windows in InstanceGiData.geo), barycentric-interpolated normals/UVs at ray hits, and textured hit albedo through a 256-slot texture binding array in its own bind group (white-padded; feature- and limit-gated with a card-albedo fallback variant). Roughness/metalness plumbed per instance for the upcoming GGX lobe. Root-cause fix: current_inv_vp_matrix is stored transposed relative to what WGSL M*v computes, so every ray the PT kernel generated since PT-1 collapsed into one degenerate bundle - transport traced garbage while the image stayed plausible (primaries come from the G-buffer). The uniform now uploads the transpose. Found via the new numeric-readback debug modes (BLOOM_PT_DEBUG=16/17 dump t/instance/prim and p0/depth to pt_trace_dump.txt); the tonemap-immune t-vs-G-buffer sanity view (debug 13) is now the mandatory first check for ray-gen changes. wgpu 29 DX12+DXC inline ray query itself verified correct with a standalone repro (fat strides, same-encoder build+dispatch, multiple queries, engine-shaped bind groups) - no naga patching needed. Tickets: docs/pt/PT-2-real-hit-shading.md (new) + PT-1 post-scriptum. --- docs/pt/PT-1-progressive-megakernel.md | 29 ++ docs/pt/PT-2-real-hit-shading.md | 55 ++++ native/shared/src/attach.rs | 12 + native/shared/src/renderer/mod.rs | 213 ++++++++++++- native/shared/src/renderer/pt_pass.rs | 157 +++++++++- native/shared/src/renderer/shaders/gi.rs | 3 + native/shared/src/renderer/shaders/pt.rs | 337 ++++++++++++++++++++- native/shared/src/renderer/shaders/ssgi.rs | 4 + native/windows/Cargo.lock | 4 +- native/windows/src/lib.rs | 22 +- 10 files changed, 807 insertions(+), 29 deletions(-) create mode 100644 docs/pt/PT-2-real-hit-shading.md diff --git a/docs/pt/PT-1-progressive-megakernel.md b/docs/pt/PT-1-progressive-megakernel.md index 13f03ae..daffd56 100644 --- a/docs/pt/PT-1-progressive-megakernel.md +++ b/docs/pt/PT-1-progressive-megakernel.md @@ -86,6 +86,35 @@ Implementation notes discovered on the way: directly. Careful with F12-capture scripts: the title screen treats the posted F12 as "press any key" and starts the game. +## Post-scriptum (2026-07-13, PT-2 bring-up): the transposed inv_vp + +PT-2's numeric readback (BLOOM_PT_DEBUG=16/17) revealed that EVERY ray +this kernel generated since PT-1 was degenerate: `current_inv_vp_matrix` +is stored transposed relative to what WGSL's `M * v` computes, so the +unprojection collapsed all primary rays into one bundle. Light transport +(sun shadow rays, bounces) traced garbage while the image still looked +plausible — primaries come from the G-buffer and direct NdotL carried +the shading. Fixed by uploading the transpose (pt_pass.rs). Lessons, +paid for in full: + +- The house rule exists for a reason: every other pass unprojects via + linear-z + projection coefficients and never touches an inverted + matrix. The PT kernel was the first `inv_vp` consumer, and the first + casualty. +- "Looks plausible" is not a transport test. The M2 convergence gate + passed on top of broken rays because accumulation mechanics are + independent of ray correctness. The debug-13 t-vs-G-buffer sanity + view (green/red/blue) is now the mandatory first check for any + ray-generation change. +- HDR-scaled probe colours (×20-50 to defeat tonemapping) saturate + fract-band signals into uniform blobs — three investigation probes + lied because of this. When probes disagree with expectations twice, + switch to numeric readback (debug 16/17) immediately. +- wgpu 29 DX12+DXC inline ray query is fully correct (validated with a + standalone repro: fat vertex strides, same-encoder build+dispatch, + multiple queries, engine-shaped bind groups — all fine). The naga + helper-function codegen needs no patching. + ## Known honest limits (by design, closed by PT-2) Flat per-mesh hit normals; card-resolution albedo at hits; no specular diff --git a/docs/pt/PT-2-real-hit-shading.md b/docs/pt/PT-2-real-hit-shading.md new file mode 100644 index 0000000..7b1d889 --- /dev/null +++ b/docs/pt/PT-2-real-hit-shading.md @@ -0,0 +1,55 @@ +# PT-2 — Real hit shading (Tier 2) + +Parent: [pt-roadmap.md](pt-roadmap.md). Builds on PT-1. + +## Scope + +Replace PT-1's flat-normal/card-albedo hit shading with real surface +data at every bounce: + +1. **Geometry megabuffer** — every TLAS instance's `Vertex3D` + index + data concatenated into two storage buffers (raw f32/u32 words, + zero repack via bytemuck; nodes retain CPU vertices). Per-instance + window in `InstanceGiData.geo` = (vertex_base, index_base, + index_count, texture_idx); `index_count == 0` falls back to PT-1 + shading. Rebuilt with the instance-data buffer (grow-only). +2. **Interpolated attributes** — `fetch_hit_attrs` reads the hit + triangle via `primitive_index`, interpolates normal + UV with + `barycentrics` (DXR convention: w=1-u-v weights v0), transforms the + normal by the ray query's `world_to_object` (v·M = (M⁻¹)ᵀ·n). +3. **Texture binding array** — `binding_array>` in its + own bind group (wgpu forbids binding arrays next to uniform + buffers), fixed 256 slots padded with the white texture (no + PARTIALLY_BOUND needed). Requires TEXTURE_BINDING_ARRAY + + SAMPLED_..._NON_UNIFORM_INDEXING **and** + `max_binding_array_elements_per_shader_stage` (defaults to 0 — + must be requested at device creation). Hit albedo = + `inst.albedo × pt_textures[geo.w]` sampled at the interpolated UV. + Without the features the kernel compiles a white-stub variant and + stays on card albedo. +4. `mat_params` (roughness/metalness) plumbed per instance — GGX + specular lobe still TODO (M3). + +## Status (2026-07-13, AMD 760M / DX12+DXC 1.9) + +- Megabuffer + windows verified end-to-end: numeric dump shows correct + per-pixel t/instance/prim; debug 6 shows smooth interpolated normals + (leaf cards each with their own direction, smooth terrain); debug 7 + shows real leaf/bark/plaster textures at traced hits. +- Debug 13 sanity: traced t agrees with the G-buffer everywhere a + proxy matches the visible mesh; expected mismatches only (grass and + skinned characters are not in the TLAS; tree proxies are unrotated). +- Found + fixed the transposed `inv_vp` upload (see the post-scriptum + in PT-1's ticket) — this had silently broken all PT-1 transport. + +## Open (M3) + +- GGX specular bounce using `mat_params` (port bloom-reference + `sample_brdf`: VNDF + Burley, Fresnel lobe pick). +- Emissive triangles at hits already work via `albedo × emissive_luma`; + emissive-as-light-source sampling is PT-4 (ReSTIR) territory. +- Known gaps, accepted: skinned enemies not in the TLAS (no per-frame + BLAS path yet); tree GI proxies unrotated (bounce-grade accuracy + only); leaf-card cutouts opaque to rays. +- Point-light NEE + colour-bleed visual + bloom-reference RMSE (carried + from PT-1 M3) — needs a lit test world + fixed exposure. diff --git a/native/shared/src/attach.rs b/native/shared/src/attach.rs index f34d8cf..66f578f 100644 --- a/native/shared/src/attach.rs +++ b/native/shared/src/attach.rs @@ -126,6 +126,13 @@ pub unsafe fn attach_engine( if !force_sw_gi && supported.contains(rt_mask) { required_features |= rt_mask; } + // PT-2: texture binding array + non-uniform indexing for textured + // path-trace hit shading (mirrors bloom_init_window). + let pt_tex_mask = wgpu::Features::TEXTURE_BINDING_ARRAY + | wgpu::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING; + if supported.contains(pt_tex_mask) { + required_features |= pt_tex_mask; + } let experimental_features = if required_features.intersects(rt_mask) { // wgpu 29 requires this explicit opt-in token for EXPERIMENTAL_* // features. Apple-Silicon Metal ray query has been stable since @@ -161,6 +168,11 @@ pub unsafe fn attach_engine( required_limits.max_samplers_per_shader_stage = required_limits .max_samplers_per_shader_stage .max(adapter_limits.max_samplers_per_shader_stage); + // PT-2: binding arrays have their own element budget, default 0. + if required_features.contains(pt_tex_mask) { + required_limits.max_binding_array_elements_per_shader_stage = + adapter_limits.max_binding_array_elements_per_shader_stage; + } if required_features.intersects(rt_mask) { required_limits = diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index a1bb1dd..8caf97c 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -359,6 +359,11 @@ struct ProbeTraceParams { wsrc_cascades: [[f32; 4]; 3], } +/// PT-2 — fixed size of the kernel's texture binding array. Real texture +/// indices at or above this clamp to 0 (white); unused tail slots are +/// padded with the white texture so no PARTIALLY_BOUND feature is needed. +const PT_MAX_TEXTURES: usize = 256; + /// Uniform for the path-tracing megakernel — WGSL mirror `PtParams` in /// shaders/pt.rs. Field-for-field, 16-byte aligned, 672 bytes. #[repr(C)] @@ -610,6 +615,15 @@ struct InstanceGiDataCpu { /// bounce on non-RT adapters (round-2 audit F4). world_aabb_min: [f32; 4], world_aabb_max: [f32; 4], + /// PT-2 — geometry window into the PT megabuffers + texture id. + /// x = first vertex (Vertex3D-stride slot) in pt_geo_vertices, + /// y = first index in pt_geo_indices, z = index count, + /// w = albedo texture index (renderer texture store; 0 = white). + /// z == 0 marks "no geometry window" (kernel falls back to the + /// flat normal + card albedo, i.e. PT-1 behaviour). + geo: [u32; 4], + /// PT-2 — x = roughness, y = metalness, z/w unused. + mat_params: [f32; 4], } #[repr(C)] @@ -1260,6 +1274,25 @@ pub struct Renderer { pt_last_tlas_version: u64, /// BLOOM_PT_DEBUG view selector, forwarded to the kernel via cfg.w. pt_debug: f32, + /// PT-2 — concatenated Vertex3D words (f32) / indices (u32) for + /// interpolated hit shading. Grow-only, rebuilt alongside the + /// instance-data buffer. + pt_geo_vertex_buffer: Option, + pt_geo_index_buffer: Option, + /// PT-2 — adapter grants binding-array features; when false the + /// kernel compiles without the texture array and hit shading stays + /// on card albedo. + pt_texture_arrays_enabled: bool, + /// PT-2 — the texture binding array lives in its own group (wgpu + /// forbids binding arrays next to uniform buffers). + pt_tex_layout: Option, + pt_tex_bg: Option, + /// Texture-store length baked into the current pt_tex_bg; growth + /// forces a rebuild so new textures become visible to PT. + pt_bg_texture_count: usize, + /// Debug-16 numeric readback (see pt_pass.rs). + pt_readback_buffer: Option, + pt_dump_written: bool, /// Ticket 014 V3 — SW SDF sphere-trace pipeline + layout. /// Active whenever the scene clipmap has been baked; chosen at @@ -1870,6 +1903,15 @@ impl Renderer { let hw_rt_enabled = device .features() .contains(wgpu::Features::EXPERIMENTAL_RAY_QUERY); + // PT-2 — textured hit shading wants a bindless-ish texture + // array. Both feature bits AND the element budget (a separate + // limit that defaults to 0 and must be granted at device + // creation), or the kernel compiles with the card-only stub. + let pt_texture_arrays_enabled = device.features().contains( + wgpu::Features::TEXTURE_BINDING_ARRAY + | wgpu::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING, + ) && device.limits().max_binding_array_elements_per_shader_stage + >= PT_MAX_TEXTURES as u32; // --- Shaders --- let shader_2d = device.create_shader_module(wgpu::ShaderModuleDescriptor { @@ -4965,15 +5007,30 @@ impl Renderer { // Same gate as the HW probe trace: no ray query, no path tracer — // and deliberately no software fallback (a CPU-speed path trace is // seconds per frame, worse than useless in any shipped mode). - let (pt_pipeline, pt_layout) = if hw_rt_enabled { - let pt_source = format!("enable wgpu_ray_query;\n{}", PT_KERNEL_WGSL); + let (pt_pipeline, pt_layout, pt_tex_layout) = if hw_rt_enabled { + // PT-2 — the texture-array half of hit shading is feature- + // gated. The variant snippet defines PT_HAS_TEXTURES plus + // pt_tex_sample; without the features the kernel compiles + // with a constant-white stub and hit shading stays on the + // card-atlas path (PT-1 behaviour). + let tex_variant = if pt_texture_arrays_enabled { + // Own bind group: wgpu forbids mixing a binding array + // with a uniform buffer inside one group. + "const PT_HAS_TEXTURES: bool = true;\n\ + @group(1) @binding(0) var pt_textures: binding_array>;\n\ + fn pt_tex_sample(idx: u32, uv: vec2) -> vec3 {\n\ + return textureSampleLevel(pt_textures[idx], card_samp, uv, 0.0).rgb;\n\ + }\n" + } else { + "const PT_HAS_TEXTURES: bool = false;\n\ + fn pt_tex_sample(idx: u32, uv: vec2) -> vec3 { return vec3(1.0); }\n" + }; + let pt_source = format!("enable wgpu_ray_query;\n{}\n{}", PT_KERNEL_WGSL, tex_variant); let pt_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("pt_kernel_shader"), source: wgpu::ShaderSource::Wgsl(pt_source.into()), }); - let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("pt_layout"), - entries: &[ + let mut pt_layout_entries = vec![ wgpu::BindGroupLayoutEntry { binding: 0, visibility: wgpu::ShaderStages::COMPUTE, ty: wgpu::BindingType::Buffer { @@ -5047,11 +5104,53 @@ impl Renderer { view_dimension: wgpu::TextureViewDimension::D2, }, count: None, }, - ], + // PT-2 — geometry megabuffers (vertex words + indices). + wgpu::BindGroupLayoutEntry { + binding: 10, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 11, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + ]; + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("pt_layout"), + entries: &pt_layout_entries, }); + // PT-2 — the texture array lives in its own group: wgpu + // forbids a binding array next to a uniform buffer. Fixed + // size; the bind group pads unused slots with the white + // texture (slot 0), so no PARTIALLY_BOUND requirement. + let tex_layout = if pt_texture_arrays_enabled { + Some(device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("pt_tex_layout"), + entries: &[wgpu::BindGroupLayoutEntry { + binding: 0, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, + count: Some(std::num::NonZeroU32::new(PT_MAX_TEXTURES as u32).unwrap()), + }], + })) + } else { + None + }; + let mut pl_groups: Vec> = vec![Some(&layout)]; + if let Some(tl) = tex_layout.as_ref() { + pl_groups.push(Some(tl)); + } let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("pt_pl_layout"), - bind_group_layouts: &[Some(&layout)], + bind_group_layouts: &pl_groups, immediate_size: 0, }); let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { @@ -5060,9 +5159,9 @@ impl Renderer { module: &pt_shader, entry_point: Some("cs_main"), compilation_options: Default::default(), cache: None, }); - (Some(pipeline), Some(layout)) + (Some(pipeline), Some(layout), tex_layout) } else { - (None, None) + (None, None, None) }; let pt_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_uniform"), @@ -6683,6 +6782,14 @@ impl Renderer { pt_prev_vp: [[0.0; 4]; 4], pt_last_tlas_version: 0, pt_debug, + pt_geo_vertex_buffer: None, + pt_geo_index_buffer: None, + pt_texture_arrays_enabled, + pt_tex_layout, + pt_tex_bg: None, + pt_bg_texture_count: 0, + pt_readback_buffer: None, + pt_dump_written: false, probe_trace_sdf_pipeline, probe_trace_sdf_layout, probe_trace_sdf_bg_cache: [None, None], @@ -7879,6 +7986,15 @@ impl Renderer { // EN-023 — running mean of instance albedos feeds the SW WSRC // bake's ground-bounce term. let mut albedo_sum = [0.0f32; 3]; + // PT-2 — geometry megabuffers: every instance's Vertex3D + index + // data concatenated so the path-trace kernel can interpolate + // real normals/UVs at ray hits. Raw Vertex3D stride (24 f32); + // instances without retained CPU geometry get index_count = 0 + // and the kernel falls back to flat-normal/card shading. + // Shared meshes are duplicated per node (nodes own their + // vertices); arena-scale worlds measure a few tens of MB. + let mut geo_vertices: Vec = Vec::new(); + let mut geo_indices: Vec = Vec::new(); for &h in instance_handles { let n = scene.nodes.get(h).unwrap(); let e = n.material.emissive; @@ -7886,6 +8002,16 @@ impl Renderer { Some(s) => (s as f32, 1.0_f32), None => (0.0, 0.0), }; + let (vertex_base, index_base, index_count) = + if !n.vertices.is_empty() && !n.indices.is_empty() { + let vb = (geo_vertices.len() / 24) as u32; + let ib = geo_indices.len() as u32; + geo_vertices.extend_from_slice(bytemuck::cast_slice(&n.vertices)); + geo_indices.extend_from_slice(&n.indices); + (vb, ib, n.indices.len() as u32) + } else { + (0, 0, 0) + }; // EN-023 — world AABB for the SDF broad-phase. The scene's // bounds pass keeps world_bounds fresh; the sentinel // (min.x > max.x = not yet computed) falls back to the local @@ -7908,8 +8034,62 @@ impl Renderer { card_aabb_max: [n.bounds_max[0], n.bounds_max[1], n.bounds_max[2], 0.0], world_aabb_min: [wmin[0], wmin[1], wmin[2], 0.0], world_aabb_max: [wmax[0], wmax[1], wmax[2], 0.0], + geo: [ + vertex_base, + index_base, + index_count, + // Indices beyond the kernel's fixed array clamp to + // white — same render as an unloaded texture. + if (n.material.texture_idx as usize) < PT_MAX_TEXTURES { + n.material.texture_idx + } else { + 0 + }, + ], + mat_params: [n.material.roughness, n.material.metalness, 0.0, 0.0], }); } + // PT-2 — upload the megabuffers (grow-only; a minimum size keeps + // bind-group creation valid before any geometry is committed). + { + let v_bytes = (geo_vertices.len() * 4).max(16) as u64; + let i_bytes = (geo_indices.len() * 4).max(16) as u64; + let v_recreate = self.pt_geo_vertex_buffer.as_ref().map_or(true, |b| b.size() < v_bytes); + let i_recreate = self.pt_geo_index_buffer.as_ref().map_or(true, |b| b.size() < i_bytes); + if v_recreate { + self.pt_geo_vertex_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_geo_vertices"), + size: v_bytes.next_power_of_two(), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + } + if i_recreate { + self.pt_geo_index_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_geo_indices"), + size: i_bytes.next_power_of_two(), + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + } + if v_recreate || i_recreate { + self.pt_bg = None; + } + if !geo_vertices.is_empty() { + self.queue.write_buffer( + self.pt_geo_vertex_buffer.as_ref().unwrap(), + 0, + bytemuck::cast_slice(&geo_vertices), + ); + } + if !geo_indices.is_empty() { + self.queue.write_buffer( + self.pt_geo_index_buffer.as_ref().unwrap(), + 0, + bytemuck::cast_slice(&geo_indices), + ); + } + } if !instance_data.is_empty() { let inv = 1.0 / instance_data.len() as f32; self.gi_scene_avg_albedo = [ @@ -7918,6 +8098,21 @@ impl Renderer { albedo_sum[2] * inv, ]; } + // PT-2 diagnostic (BLOOM_PT_DEBUG only): eprintln is unreliable + // after init on Windows, so drop a summary file in cwd. + if self.pt_debug > 0.0 { + let with_geo = instance_data.iter().filter(|d| d.geo[2] > 0).count(); + let _ = std::fs::write( + "pt_geo_debug.txt", + format!( + "instances={} with_geo_window={} megabuffer_verts={} megabuffer_indices={}\n", + instance_data.len(), + with_geo, + geo_vertices.len() / 24, + geo_indices.len(), + ), + ); + } if !instance_data.is_empty() { self.queue.write_buffer( self.tlas_instance_data_buffer.as_ref().unwrap(), diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index 4b590f6..a8d9ce5 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -25,10 +25,17 @@ impl Renderer { if self.pt_pipeline.is_none() || self.tlas.is_none() || self.tlas_instance_data_buffer.is_none() + || self.pt_geo_vertex_buffer.is_none() + || self.pt_geo_index_buffer.is_none() { self.pt_accum_count = 0; return; } + // PT-2 — a grown texture store means the baked view array is + // stale; rebuild so new textures become visible to hit shading. + if self.pt_texture_arrays_enabled && self.pt_bg_texture_count != self.textures.len() { + self.pt_tex_bg = None; + } // ---- accumulation validity ---- // Any camera motion beyond epsilon restarts progressive @@ -75,7 +82,11 @@ impl Renderer { self.pt_accum_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_accum"), size: needed, - usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + // COPY_SRC: the debug-16 numeric readback copies a window + // of this buffer to a staging buffer. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false, })); self.pt_bg = None; @@ -102,8 +113,22 @@ impl Renderer { let max_bounces = if self.pt_mode == 2 { 2.0 } else { 8.0 }; let cam = self.current_camera_pos; + // current_inv_vp_matrix is stored transposed relative to what + // WGSL's `M * v` needs (the composed VP inherits mat4_multiply's + // convention; its inverse lands transposed). Upload the + // transpose so the kernel's unprojection is the real inverse — + // without this every ray collapses to one degenerate bundle and + // the whole path trace silently hits garbage (found via numeric + // readback; see docs/pt/PT-2 notes). + let m = &self.current_inv_vp_matrix; + let inv_vp_t = [ + [m[0][0], m[1][0], m[2][0], m[3][0]], + [m[0][1], m[1][1], m[2][1], m[3][1]], + [m[0][2], m[1][2], m[2][2], m[3][2]], + [m[0][3], m[1][3], m[2][3], m[3][3]], + ]; let params = PtParamsCpu { - inv_vp: self.current_inv_vp_matrix, + inv_vp: inv_vp_t, cam_pos: [cam[0], cam[1], cam[2], 0.0], sun_dir: [ -ld[0] * sun_inv_len, @@ -138,10 +163,7 @@ impl Renderer { // buffer recreation) ---- if self.pt_bg.is_none() { let tlas = self.tlas.as_ref().unwrap(); - self.pt_bg = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("pt_bg"), - layout: self.pt_layout.as_ref().unwrap(), - entries: &[ + let entries = vec![ wgpu::BindGroupEntry { binding: 0, resource: self.pt_uniform_buffer.as_entire_binding() }, wgpu::BindGroupEntry { binding: 1, resource: tlas.as_binding() }, wgpu::BindGroupEntry { binding: 2, resource: self.tlas_instance_data_buffer.as_ref().unwrap().as_entire_binding() }, @@ -155,8 +177,36 @@ impl Renderer { wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::Sampler(&self.mesh_card_atlas_sampler) }, wgpu::BindGroupEntry { binding: 8, resource: self.pt_accum_buffer.as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, - ], + wgpu::BindGroupEntry { binding: 10, resource: self.pt_geo_vertex_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 11, resource: self.pt_geo_index_buffer.as_ref().unwrap().as_entire_binding() }, + ]; + self.pt_bg = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_bg"), + layout: self.pt_layout.as_ref().unwrap(), + entries: &entries, + })); + } + // PT-2 — group 1: the texture binding array. Real store views + // first, white (slot 0) padding to the fixed layout count. The + // bind group holds refs, so the temporary views live with it. + if self.pt_texture_arrays_enabled && self.pt_tex_bg.is_none() { + let n = self.textures.len().min(PT_MAX_TEXTURES); + let tex_views: Vec = (0..n.max(1)) + .map(|i| self.textures[i.min(self.textures.len() - 1)] + .create_view(&wgpu::TextureViewDescriptor::default())) + .collect(); + let tex_view_refs: Vec<&wgpu::TextureView> = (0..PT_MAX_TEXTURES) + .map(|i| &tex_views[if i < n { i } else { 0 }]) + .collect(); + self.pt_tex_bg = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_tex_bg"), + layout: self.pt_tex_layout.as_ref().unwrap(), + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureViewArray(&tex_view_refs), + }], })); + self.pt_bg_texture_count = self.textures.len(); } // ---- dispatch ---- @@ -168,8 +218,101 @@ impl Renderer { }); pass.set_pipeline(self.pt_pipeline.as_ref().unwrap()); pass.set_bind_group(0, self.pt_bg.as_ref().unwrap(), &[]); + if self.pt_texture_arrays_enabled { + pass.set_bind_group(1, self.pt_tex_bg.as_ref().unwrap(), &[]); + } pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); } self.pt_accum_count = self.pt_accum_count.saturating_add(1); + + // ---- debug 16: numeric readback of traced intersections ---- + // Copies a window of the accum buffer (center of frame) into a + // staging buffer each frame; the previous frame's copy is mapped + // (blocking) and dumped to pt_trace_dump.txt once. + if (self.pt_debug == 16.0 || self.pt_debug == 17.0) + && self.pt_accum_count > 30 + && !self.pt_dump_written + { + let dump_pixels: u64 = (surf_w as u64).min(4096); + let row = (surf_h / 2) as u64; + let offset = row * surf_w as u64 * 16; + if self.pt_readback_buffer.is_none() { + self.pt_readback_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_readback"), + size: dump_pixels * 16, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + encoder.copy_buffer_to_buffer( + self.pt_accum_buffer.as_ref().unwrap(), + offset, + self.pt_readback_buffer.as_ref().unwrap(), + 0, + dump_pixels * 16, + ); + } else { + // Previous frame's copy has been submitted; map it now. + let buf = self.pt_readback_buffer.as_ref().unwrap(); + let slice = buf.slice(..); + slice.map_async(wgpu::MapMode::Read, |_| {}); + let _ = self.device.poll(wgpu::PollType::Wait { submission_index: None, timeout: None }); + let data = slice.get_mapped_range(); + let vals: &[[f32; 4]] = bytemuck::cast_slice(&data); + let mut out = String::new(); + out.push_str(&format!( + "middle row, {} pixels, mode {}\n", + vals.len(), + self.pt_debug + )); + // Every 64th pixel across the full row. Field meaning: + // 16 = t / id / prim / kind; 17 = p0.xyz / raw depth. + for (i, v) in vals.iter().enumerate().step_by(64) { + out.push_str(&format!( + "col {i}: {:.4} {:.4} {:.4} {:.6}\n", + v[0], v[1], v[2], v[3] + )); + } + // Also dump the CPU-side uniform inputs for comparison, + // plus the unprojection computed in BOTH multiply + // conventions. Whichever matches the GPU dump is what + // the shader effectively computed; the other (if sane) + // is the fix. + let ndc = [0.0f32, 0.0, 0.998647, 1.0]; + let m = &self.current_inv_vp_matrix; + let mut h_col = [0.0f32; 4]; // h_i = sum_c m[c][i] * ndc[c] + let mut h_row = [0.0f32; 4]; // h_i = sum_c m[i][c] * ndc[c] + for i in 0..4 { + for c in 0..4 { + h_col[i] += m[c][i] * ndc[c]; + h_row[i] += m[i][c] * ndc[c]; + } + } + out.push_str(&format!( + "cpu cam_pos = {:?}\n\ + unproject as columns: h={:?} p={:?}\n\ + unproject transposed: h={:?} p={:?}\n", + self.current_camera_pos, + h_col, + [h_col[0] / h_col[3], h_col[1] / h_col[3], h_col[2] / h_col[3]], + h_row, + [h_row[0] / h_row[3], h_row[1] / h_row[3], h_row[2] / h_row[3]], + )); + // Distinct instance-id count over the whole row. + let mut ids: Vec = vals + .iter() + .filter(|v| v[3] != 0.0) + .map(|v| v[1] as i64) + .collect(); + ids.sort_unstable(); + ids.dedup(); + let misses = vals.iter().filter(|v| v[3] == 0.0).count(); + out.push_str(&format!("distinct hit ids: {:?}\n", ids)); + out.push_str(&format!("misses: {}\n", misses)); + drop(data); + buf.unmap(); + let _ = std::fs::write("pt_trace_dump.txt", out); + self.pt_dump_written = true; + } + } } } diff --git a/native/shared/src/renderer/shaders/gi.rs b/native/shared/src/renderer/shaders/gi.rs index 9b228c4..8f81dae 100644 --- a/native/shared/src/renderer/shaders/gi.rs +++ b/native/shared/src/renderer/shaders/gi.rs @@ -621,6 +621,9 @@ struct HwBakeInstanceGiData { // EN-023 — world-space AABB (SDF path only; layout mirror). world_aabb_min: vec4, world_aabb_max: vec4, + // PT-2 — layout mirror only; the WSRC bake ignores both fields. + geo: vec4, + mat_params: vec4, }; const HW_BAKE_CARD_SLOTS_PER_ROW: f32 = 64.0; diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index ad1b977..535d41b 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -24,6 +24,20 @@ //! 1 = raw depth visualised 2 = reconstructed world normals //! 3 = G-buffer albedo 4 = sun shadow-ray visibility //! 5 = solid magenta (pipeline probe — proves dispatch + write path) +//! 6 = traced-primary interpolated normal (compare against 2; +//! magenta = TLAS miss where G-buffer had geometry, orange = hit +//! instance without a geometry window) +//! 7 = traced-primary textured hit albedo (compare against 3; +//! yellow = adapter lacks texture-array features) +//! 8-15 = binary/quantized probes from the DX12 bring-up (hit-window +//! flag, normal axes, primitive/instance banding, two-query +//! aliasing, t-vs-G-buffer sanity, t contours). 13 is the +//! keeper: green = traced t agrees with the G-buffer, red = +//! mismatch, blue = miss. +//! 16/17 = NUMERIC dumps via the accum buffer + CPU readback +//! (pt_trace_dump.txt): 16 = t/instance/prim/kind, 17 = p0 + +//! raw depth. These found the transposed inv_vp: when every +//! probe looks "constant", dump numbers before theorizing. pub(in crate::renderer) const PT_KERNEL_WGSL: &str = r#" struct PtLight { @@ -53,6 +67,11 @@ struct InstanceGiData { card_aabb_max: vec4, world_aabb_min: vec4, world_aabb_max: vec4, + // PT-2: x = vertex_base, y = index_base, z = index_count (0 = no + // geometry window -> PT-1 fallback), w = albedo texture index. + geo: vec4, + // PT-2: x = roughness, y = metalness. + mat_params: vec4, }; @group(0) @binding(0) var u: PtParams; @@ -65,6 +84,56 @@ struct InstanceGiData { @group(0) @binding(7) var card_samp: sampler; @group(0) @binding(8) var accum: array>; @group(0) @binding(9) var out_hdr: texture_storage_2d; +// PT-2: geometry megabuffers. geo_v holds raw Vertex3D words (stride 24 +// f32: position +0, normal +3, color +6, uv +10, ...); geo_i holds the +// concatenated index streams. Windows are per-instance via inst.geo. +// (Binding 12, the texture array + PT_HAS_TEXTURES + pt_tex_sample, is +// appended by the Rust side per adapter support.) +@group(0) @binding(10) var geo_v: array; +@group(0) @binding(11) var geo_i: array; + +const PT_VSTRIDE: u32 = 24u; + +struct HitAttrs { + normal_os: vec3, + uv: vec2, +}; + +fn vert_normal_os(slot: u32) -> vec3 { + let o = slot * PT_VSTRIDE + 3u; + return vec3(geo_v[o], geo_v[o + 1u], geo_v[o + 2u]); +} + +fn vert_uv(slot: u32) -> vec2 { + let o = slot * PT_VSTRIDE + 10u; + return vec2(geo_v[o], geo_v[o + 1u]); +} + +// Interpolate the hit triangle's vertex normal + UV. DXR/Vulkan +// barycentric convention: (u, v) weight vertices 1 and 2, w = 1-u-v +// weights vertex 0. +fn fetch_hit_attrs(geo: vec4, prim: u32, bary: vec2) -> HitAttrs { + let base = geo.y + prim * 3u; + let s0 = geo.x + geo_i[base]; + let s1 = geo.x + geo_i[base + 1u]; + let s2 = geo.x + geo_i[base + 2u]; + let w = 1.0 - bary.x - bary.y; + var a: HitAttrs; + a.normal_os = w * vert_normal_os(s0) + bary.x * vert_normal_os(s1) + bary.y * vert_normal_os(s2); + a.uv = w * vert_uv(s0) + bary.x * vert_uv(s1) + bary.y * vert_uv(s2); + return a; +} + +// Object-space normal -> world space: with M = object_to_world the +// correct transform is (M^-1)^T, and the ray query hands us M^-1 as +// world_to_object. `v * mat3` multiplies by the transpose in WGSL. +fn normal_to_world(n_os: vec3, w2o: mat4x3) -> vec3 { + let lin = mat3x3(w2o[0], w2o[1], w2o[2]); + let n = n_os * lin; + let len = length(n); + if (len < 1e-8) { return vec3(0.0, 1.0, 0.0); } + return n / len; +} // ---- RNG: PCG, one stream per (pixel, frame) -------------------------------- @@ -327,6 +396,242 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { textureStore(out_hdr, px, vec4(vec3(vis), 1.0)); return; } + if (debug == 8.0) { + // Binary probe: white = traced hit has a geometry window, + // black = geo.z reads 0, red = TLAS miss. HDR-large values so + // exposure/tonemap can't blur the verdict. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq8: ray_query; + rayQueryInitialize(&rq8, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq8)) { break; } + } + let h8 = rayQueryGetCommittedIntersection(&rq8); + var c8 = vec3(100.0, 0.0, 0.0); + if (h8.kind != RAY_QUERY_INTERSECTION_NONE) { + let gi = instance_data[h8.instance_custom_data].geo; + c8 = select(vec3(0.0), vec3(100.0), gi.z > 0u); + } + textureStore(out_hdr, px, vec4(c8, 1.0)); + return; + } + if (debug == 9.0) { + // Quantized normal probe: dominant axis of the interpolated + // world normal as six saturated HDR colours. +X red, -X dark + // red-ish magenta, +Y green, -Y cyan, +Z blue, -Z yellow. + // Gray = TLAS miss / no window / zero-length normal. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq9: ray_query; + rayQueryInitialize(&rq9, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq9)) { break; } + } + let h9 = rayQueryGetCommittedIntersection(&rq9); + var c9 = vec3(5.0, 5.0, 5.0); + if (h9.kind != RAY_QUERY_INTERSECTION_NONE) { + let inst9 = instance_data[h9.instance_custom_data]; + if (inst9.geo.z > 0u) { + let a9 = fetch_hit_attrs(inst9.geo, h9.primitive_index, h9.barycentrics); + let raw = a9.normal_os; + if (length(raw) > 1e-6) { + let n9 = normal_to_world(raw, h9.world_to_object); + let an = abs(n9); + if (an.y >= an.x && an.y >= an.z) { + c9 = select(vec3(0.0, 50.0, 50.0), vec3(0.0, 50.0, 0.0), n9.y >= 0.0); + } else if (an.x >= an.z) { + c9 = select(vec3(50.0, 0.0, 25.0), vec3(50.0, 0.0, 0.0), n9.x >= 0.0); + } else { + c9 = select(vec3(50.0, 50.0, 0.0), vec3(0.0, 0.0, 50.0), n9.z >= 0.0); + } + } + } + } + textureStore(out_hdr, px, vec4(c9, 1.0)); + return; + } + if (debug == 10.0) { + // primitive_index sanity probe: banded pseudo-colour of the hit + // triangle index. Expected: per-triangle colour noise across + // meshes. A single flat colour everywhere = the field is + // constant; saturated white = garbage-huge. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq10: ray_query; + rayQueryInitialize(&rq10, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq10)) { break; } + } + let h10 = rayQueryGetCommittedIntersection(&rq10); + var c10 = vec3(0.0); + if (h10.kind != RAY_QUERY_INTERSECTION_NONE) { + let prim = f32(h10.primitive_index); + c10 = vec3(fract(prim / 64.0), fract(prim / 1024.0), fract(prim / 16384.0)) * 30.0; + } + textureStore(out_hdr, px, vec4(c10, 1.0)); + return; + } + if (debug == 11.0 || debug == 12.0) { + // 11: instance_custom_data palette (expect distinct colours per + // proxy: terrain vs trees vs building). Constant = broken. + // 12: raw barycentrics (expect smooth per-triangle gradients). + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq11: ray_query; + rayQueryInitialize(&rq11, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq11)) { break; } + } + let h11 = rayQueryGetCommittedIntersection(&rq11); + var c11 = vec3(0.0); + if (h11.kind != RAY_QUERY_INTERSECTION_NONE) { + if (debug == 11.0) { + let id = h11.instance_custom_data; + c11 = vec3( + f32((id * 37u) % 7u) / 7.0, + f32((id * 61u) % 11u) / 11.0, + f32((id * 13u) % 5u) / 5.0, + ) * 30.0; + } else { + let b = h11.barycentrics; + c11 = vec3(b.x, b.y, max(0.0, 1.0 - b.x - b.y)) * 30.0; + } + } + textureStore(out_hdr, px, vec4(c11, 1.0)); + return; + } + if (debug == 13.0) { + // TLAS sanity: green = traced primary hit distance agrees with + // the G-buffer depth (within 2% + 0.1m), red = disagreement + // (wrong geometry committed), blue = TLAS miss on a G-buffer + // pixel. If this is red/blue everywhere the TLAS itself (not + // the intersection attributes) is broken on this backend. + let to_p = p0 - u.cam_pos.xyz; + let gdist = length(to_p); + let dir0 = to_p / max(gdist, 1e-4); + var rq13: ray_query; + rayQueryInitialize(&rq13, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq13)) { break; } + } + let h13 = rayQueryGetCommittedIntersection(&rq13); + var c13 = vec3(0.0, 0.0, 50.0); + if (h13.kind != RAY_QUERY_INTERSECTION_NONE) { + let err = abs(h13.t - gdist); + if (err < gdist * 0.02 + 0.1) { + c13 = vec3(0.0, 50.0, 0.0); + } else { + c13 = vec3(50.0, 0.0, 0.0); + } + } + textureStore(out_hdr, px, vec4(c13, 1.0)); + return; + } + if (debug == 14.0) { + // Shape probe: contour bands of the traced primary hit distance + // — shows what world the TLAS actually contains. Blue = miss. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq14: ray_query; + rayQueryInitialize(&rq14, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq14)) { break; } + } + let h14 = rayQueryGetCommittedIntersection(&rq14); + var c14 = vec3(0.0, 0.0, 30.0); + if (h14.kind != RAY_QUERY_INTERSECTION_NONE) { + c14 = vec3( + fract(h14.t * 0.125), + fract(h14.t * 0.03125), + fract(h14.t * 0.0078125), + ) * 20.0; + } + textureStore(out_hdr, px, vec4(c14, 1.0)); + return; + } + if (debug == 15.0) { + // Aliasing probe: two queries, two very different rays. + // A = primary (per-pixel), B = straight down (t ~= camera + // height, near-constant). R channel = banded tA, G = banded tB. + // If R == G everywhere the two queries alias to one object. + let dirA = normalize(p0 - u.cam_pos.xyz); + var rqA: ray_query; + rayQueryInitialize(&rqA, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dirA)); + loop { + if (!rayQueryProceed(&rqA)) { break; } + } + var rqB: ray_query; + rayQueryInitialize(&rqB, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, vec3(0.0, -1.0, 0.0))); + loop { + if (!rayQueryProceed(&rqB)) { break; } + } + let hA = rayQueryGetCommittedIntersection(&rqA); + let hB = rayQueryGetCommittedIntersection(&rqB); + var tA = -1.0; + var tB = -1.0; + if (hA.kind != RAY_QUERY_INTERSECTION_NONE) { tA = hA.t; } + if (hB.kind != RAY_QUERY_INTERSECTION_NONE) { tB = hB.t; } + let c15 = vec3(fract(tA * 0.125) * 20.0, fract(tB * 0.125) * 20.0, 0.0); + textureStore(out_hdr, px, vec4(c15, 1.0)); + return; + } + if (debug == 16.0) { + // Raw numeric dump: traced primary intersection into the accum + // buffer as (t, instance_custom_data, primitive_index, kind). + // The CPU side reads a window of this buffer back and writes a + // text file — no tonemap guesswork. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq16: ray_query; + rayQueryInitialize(&rq16, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq16)) { break; } + } + let h16 = rayQueryGetCommittedIntersection(&rq16); + let idx16 = gid.y * u.size.x + gid.x; + accum[idx16] = vec4( + h16.t, + f32(h16.instance_custom_data), + f32(h16.primitive_index), + f32(h16.kind), + ); + textureStore(out_hdr, px, vec4(0.2, 0.0, 0.4, 1.0)); + return; + } + if (debug == 17.0) { + // Raw ray-generation dump: reconstructed world position + raw + // depth, straight into accum for CPU readback. + let idx17 = gid.y * u.size.x + gid.x; + accum[idx17] = vec4(p0, depth); + textureStore(out_hdr, px, vec4(0.4, 0.2, 0.0, 1.0)); + return; + } + if (debug == 6.0 || debug == 7.0) { + // PT-2 validation: trace the primary ray through the TLAS + // (ignoring the G-buffer) and show interpolated attributes. + // Should match debug 2/3 up to smooth-vs-screen normals and + // card-vs-texture resolution. + let dir0 = normalize(p0 - u.cam_pos.xyz); + var rq0: ray_query; + rayQueryInitialize(&rq0, accel, RayDesc(0u, 0xFFu, 0.001, 1000.0, u.cam_pos.xyz, dir0)); + loop { + if (!rayQueryProceed(&rq0)) { break; } + } + let h = rayQueryGetCommittedIntersection(&rq0); + var col = vec3(1.0, 0.0, 1.0); // magenta: TLAS miss + if (h.kind != RAY_QUERY_INTERSECTION_NONE) { + let hinst = instance_data[h.instance_custom_data]; + if (hinst.geo.z > 0u) { + let attrs = fetch_hit_attrs(hinst.geo, h.primitive_index, h.barycentrics); + if (debug == 6.0) { + col = normal_to_world(attrs.normal_os, h.world_to_object) * 0.5 + vec3(0.5); + } else if (PT_HAS_TEXTURES) { + col = hinst.albedo * pt_tex_sample(hinst.geo.w, attrs.uv); + } else { + col = vec3(1.0, 1.0, 0.0); // yellow: no tex arrays + } + } else { + col = vec3(1.0, 0.5, 0.0); // orange: no geo window + } + } + textureStore(out_hdr, px, vec4(col, 1.0)); + return; + } // ---- one path sample -------------------------------------------------- @@ -354,18 +659,30 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let inst = instance_data[hit.instance_custom_data]; let hit_ws = origin + dir * hit.t; let hit_os = (hit.world_to_object * vec4(hit_ws, 1.0)).xyz; - let alb_hit = albedo_at_hit(inst, hit_os, dir); - - // Hit normal: the flat per-instance normal (Tier-1 honesty note in - // the roadmap; PT-2 replaces this with interpolated vertex normals). - var n_hit = inst.normal_ws; - let n_len = length(n_hit); - if (n_len < 1e-4) { - n_hit = -dir; + + // PT-2: interpolated vertex normal + textured albedo when the + // instance carries a geometry window; PT-1 flat-normal/card + // fallback otherwise. + var n_hit: vec3; + var alb_hit: vec3; + if (inst.geo.z > 0u) { + let attrs = fetch_hit_attrs(inst.geo, hit.primitive_index, hit.barycentrics); + n_hit = normal_to_world(attrs.normal_os, hit.world_to_object); + if (PT_HAS_TEXTURES) { + alb_hit = inst.albedo * pt_tex_sample(inst.geo.w, attrs.uv); + } else { + alb_hit = albedo_at_hit(inst, hit_os, dir); + } } else { - n_hit = n_hit / n_len; - if (dot(n_hit, dir) > 0.0) { n_hit = -n_hit; } + var nf = inst.normal_ws; + let n_len = length(nf); + if (n_len < 1e-4) { nf = -dir; } else { nf = nf / n_len; } + n_hit = nf; + alb_hit = albedo_at_hit(inst, hit_os, dir); } + // A backface (or a flat normal pointing away) still bounces + // outward, matching the OPAQUE two-sided raster convention. + if (dot(n_hit, dir) > 0.0) { n_hit = -n_hit; } // Emissive surfaces radiate; matches the Lumen fallback semantics // (albedo * emissive_luma). diff --git a/native/shared/src/renderer/shaders/ssgi.rs b/native/shared/src/renderer/shaders/ssgi.rs index bac4030..d1a3fe3 100644 --- a/native/shared/src/renderer/shaders/ssgi.rs +++ b/native/shared/src/renderer/shaders/ssgi.rs @@ -405,6 +405,10 @@ struct InstanceGiData { // EN-023 — world-space AABB (SDF path only; layout mirror). world_aabb_min: vec4, world_aabb_max: vec4, + // PT-2 — layout mirror only (path-tracer geometry window + + // material params); the GI traces ignore both fields. + geo: vec4, + mat_params: vec4, }; const CARD_SLOTS_PER_ROW: f32 = 64.0; diff --git a/native/windows/Cargo.lock b/native/windows/Cargo.lock index 238d84c..30d947a 100644 --- a/native/windows/Cargo.lock +++ b/native/windows/Cargo.lock @@ -741,9 +741,9 @@ dependencies = [ [[package]] name = "naga" -version = "29.0.1" +version = "29.0.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "aa2630921705b9b01dcdd0b6864b9562ca3c1951eecd0f0c4f5f04f61e412647" +checksum = "b2bf919621e7975acb27d881bae2fb993e0d45c8e0446e85e6272971e00dc8df" dependencies = [ "arrayvec", "bit-set", diff --git a/native/windows/src/lib.rs b/native/windows/src/lib.rs index 7a42feb..dc43a9a 100644 --- a/native/windows/src/lib.rs +++ b/native/windows/src/lib.rs @@ -651,11 +651,15 @@ unsafe fn init_engine_for_hwnd( // GI, timestamps for the profiler) are available on it. let info = adapter.get_info(); eprintln!( - "bloom: adapter '{}' ({:?}), ray_query={}, timestamps={}", + "bloom: adapter '{}' ({:?}), ray_query={}, timestamps={}, tex_arrays={}", info.name, info.backend, adapter.features().contains(wgpu::Features::EXPERIMENTAL_RAY_QUERY), adapter.features().contains(wgpu::Features::TIMESTAMP_QUERY), + adapter.features().contains( + wgpu::Features::TEXTURE_BINDING_ARRAY + | wgpu::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING + ), ); } @@ -680,6 +684,14 @@ unsafe fn init_engine_for_hwnd( if !force_sw_gi && supported.contains(rt_mask) { required_features |= rt_mask; } + // PT-2: texture binding array + non-uniform indexing for textured + // path-trace hit shading. Both or neither (the kernel indexes the + // array with a per-thread material id). + let pt_tex_mask = wgpu::Features::TEXTURE_BINDING_ARRAY + | wgpu::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING; + if supported.contains(pt_tex_mask) { + required_features |= pt_tex_mask; + } let experimental_features = if required_features.intersects(rt_mask) { unsafe { wgpu::ExperimentalFeatures::enabled() } } else { @@ -703,6 +715,14 @@ unsafe fn init_engine_for_hwnd( adapter_limits.max_sampled_textures_per_shader_stage; required_limits.max_samplers_per_shader_stage = adapter_limits.max_samplers_per_shader_stage; + // PT-2: binding arrays have their own element budget, default 0. + // Take whatever the adapter offers; the renderer checks the + // granted value against its fixed array size before compiling + // the textured kernel variant. + if required_features.contains(pt_tex_mask) { + required_limits.max_binding_array_elements_per_shader_stage = + adapter_limits.max_binding_array_elements_per_shader_stage; + } if required_features.intersects(rt_mask) { required_limits = required_limits .using_minimum_supported_acceleration_structure_values(); From b146b6d77972df21c156518a4f8b9c5a37477b62 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 18:31:38 +0200 Subject: [PATCH 04/16] feat(pt): PT-2 GGX specular + Burley diffuse BRDF sampling Port of bloom-reference sample_brdf: Heitz VNDF specular lobe + cosine-sampled Burley diffuse, lobe picked by Fresnel-at-normal weight, metal/dielectric split via f0 = mix(0.04, base_color, metal). Primary hits read metallic/roughness from the material G-buffer; bounce hits read the per-instance mat_params plumbed by the megabuffer commit. Diffuse-only NEE is scaled by (1 - metallic); specular NEE is a documented gap (spec reflections of sky/scene come from the sampled bounce itself). Physical-convention throughput (BRDF*cos/pdf) collapses to plain albedo for pure diffuse, so the pi-premultiplied light intensities and PT-1 brightness parity are unchanged. --- native/shared/src/renderer/shaders/pt.rs | 142 ++++++++++++++++++++++- 1 file changed, 137 insertions(+), 5 deletions(-) diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 535d41b..bd9e212 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -246,6 +246,120 @@ fn sky_radiance(dir: vec3) -> vec3 { return u.sky_color.rgb * mix(0.45, 1.35, t); } +// ---- GGX BRDF sampling (PT-2; port of bloom-reference sample_brdf) -------- + +fn fresnel_schlick3(cos_theta: f32, f0: vec3) -> vec3 { + let m = clamp(1.0 - cos_theta, 0.0, 1.0); + let m2 = m * m; + return f0 + (vec3(1.0) - f0) * (m2 * m2 * m); +} + +fn smith_g1(n_dot_x: f32, alpha: f32) -> f32 { + let a2 = alpha * alpha; + let inner = sqrt((1.0 - a2) * n_dot_x * n_dot_x + a2); + return 2.0 * n_dot_x / (n_dot_x + inner + 1e-6); +} + +fn v_smith(n_dot_v: f32, n_dot_l: f32, alpha: f32) -> f32 { + let a2 = alpha * alpha; + let ggx_v = n_dot_l * sqrt((n_dot_v * (1.0 - a2) + a2) * n_dot_v); + let ggx_l = n_dot_v * sqrt((n_dot_l * (1.0 - a2) + a2) * n_dot_l); + return 0.5 / (ggx_v + ggx_l + 1e-6); +} + +fn burley_diffuse(n_dot_l: f32, n_dot_v: f32, l_dot_h: f32, roughness: f32) -> f32 { + let fd90 = 0.5 + 2.0 * l_dot_h * l_dot_h * roughness; + let ml = pow(1.0 - n_dot_l, 5.0); + let mv = pow(1.0 - n_dot_v, 5.0); + return (1.0 + (fd90 - 1.0) * ml) * (1.0 + (fd90 - 1.0) * mv) / 3.14159265; +} + +// Heitz 2018 VNDF sampler — visible-normal distribution, tangent frame. +fn sample_ggx_vndf(v_t: vec3, alpha: f32, r2: vec2) -> vec3 { + let vh = normalize(vec3(alpha * v_t.x, alpha * v_t.y, v_t.z)); + let lensq = vh.x * vh.x + vh.y * vh.y; + var t1 = vec3(1.0, 0.0, 0.0); + if (lensq > 0.0) { + t1 = vec3(-vh.y, vh.x, 0.0) / sqrt(lensq); + } + let t2 = cross(vh, t1); + let r = sqrt(r2.x); + let phi = 6.2831853 * r2.y; + let t1v = r * cos(phi); + var t2v = r * sin(phi); + let s = 0.5 * (1.0 + vh.z); + t2v = (1.0 - s) * sqrt(max(0.0, 1.0 - t1v * t1v)) + s * t2v; + let nh = t1v * t1 + t2v * t2 + sqrt(max(0.0, 1.0 - t1v * t1v - t2v * t2v)) * vh; + return normalize(vec3(alpha * nh.x, alpha * nh.y, max(nh.z, 0.0))); +} + +struct BrdfSample { + dir: vec3, + // BRDF * cos / pdf, physical convention. For the pure-diffuse case + // this reduces to plain albedo, so the game's pi-premultiplied + // light intensities are unaffected. + weight: vec3, + valid: bool, +}; + +fn sample_brdf( + n: vec3, + view_ws: vec3, + base_color: vec3, + roughness: f32, + metallic: f32, +) -> BrdfSample { + var out: BrdfSample; + out.valid = false; + let alpha = max(roughness * roughness, 1e-3); + let m = onb(n); // columns (t, bt, n): local -> world + let v_t = vec3(dot(view_ws, m[0]), dot(view_ws, m[1]), dot(view_ws, n)); + if (v_t.z <= 0.0) { + return out; + } + let f0 = mix(vec3(0.04), base_color, metallic); + let spec_weight = (f0.x + f0.y + f0.z) / 3.0; + let diff_weight = (1.0 - spec_weight) * (1.0 - metallic); + let p_spec = spec_weight / (spec_weight + diff_weight + 1e-6); + let r2 = rand_2f(); + if (rand_f() < p_spec) { + let h_t = sample_ggx_vndf(v_t, alpha, r2); + let l_t = reflect(-v_t, h_t); + if (l_t.z <= 0.0) { + return out; + } + let n_dot_l = l_t.z; + let n_dot_v = max(v_t.z, 1e-4); + let v_dot_h = max(dot(v_t, h_t), 1e-4); + let f = fresnel_schlick3(v_dot_h, f0); + // VNDF pdf: throughput collapses to F * G2 / G1(V). + let g2 = v_smith(n_dot_v, n_dot_l, alpha) * 4.0 * n_dot_v * n_dot_l; + let g1_v = smith_g1(n_dot_v, alpha); + out.dir = m * l_t; + out.weight = f * g2 / (max(g1_v, 1e-6) * max(p_spec, 1e-6)); + out.valid = true; + return out; + } + // Diffuse lobe: cosine hemisphere; weight = albedo * burley * pi + // (Burley divides by pi internally; pdf = cos/pi cancels the cos). + let r = sqrt(r2.x); + let phi = 6.2831853 * r2.y; + let l_t = vec3(r * cos(phi), r * sin(phi), sqrt(max(0.0, 1.0 - r2.x))); + let n_dot_l = max(l_t.z, 1e-4); + let n_dot_v = max(v_t.z, 1e-4); + let h_un = v_t + l_t; + var l_dot_h = 0.0; + if (dot(h_un, h_un) > 1e-8) { + l_dot_h = max(dot(l_t, normalize(h_un)), 0.0); + } + let diffuse_albedo = base_color * (1.0 - metallic) * (vec3(1.0) - f0); + let fd = burley_diffuse(n_dot_l, n_dot_v, l_dot_h, roughness); + out.dir = m * l_t; + out.weight = diffuse_albedo * fd * 3.14159265 / max(1.0 - p_spec, 1e-6); + out.valid = true; + return out; +} + // ---- Ray casts ------------------------------------------------------------------ fn occluded(origin: vec3, dir: vec3, max_t: f32) -> bool { @@ -635,14 +749,29 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // ---- one path sample -------------------------------------------------- - var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0); - var throughput = albedo0; + // Primary surface material from the G-buffer (R = metallic, + // G = roughness). NEE stays diffuse-only, so scale it by + // (1 - metallic) — metals have no diffuse lobe. Specular NEE is a + // known gap (see the PT-2 ticket); specular reflection of sky and + // scene comes from the GGX bounce below. + let mr0 = textureLoad(material_tex, px, 0).rg; + var metal_cur = mr0.r; + var rough_cur = mr0.g; + var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur)); + var throughput = vec3(1.0); var origin = p0 + n0 * 0.02; var n_cur = n0; + var alb_cur = albedo0; + var view_cur = normalize(u.cam_pos.xyz - p0); let max_bounces = u32(u.cfg.y); for (var b = 0u; b < max_bounces; b = b + 1u) { - let dir = cosine_sample(n_cur, rand_2f()); + let s = sample_brdf(n_cur, view_cur, alb_cur, rough_cur, metal_cur); + if (!s.valid) { + break; + } + throughput *= s.weight; + let dir = s.dir; var rq: ray_query; rayQueryInitialize(&rq, accel, RayDesc(0u, 0xFFu, 0.001, 500.0, origin, dir)); @@ -689,11 +818,14 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { radiance += throughput * inst.albedo * inst.emissive_luma; let hit_p = hit_ws + n_hit * 0.02; - radiance += throughput * direct_light(hit_p, n_hit, alb_hit); + radiance += throughput * direct_light(hit_p, n_hit, alb_hit * (1.0 - inst.mat_params.y)); - throughput *= alb_hit; origin = hit_p; n_cur = n_hit; + alb_cur = alb_hit; + rough_cur = inst.mat_params.x; + metal_cur = inst.mat_params.y; + view_cur = -dir; // Russian roulette from the third bounce. if (b >= 2u) { From dfbc54aedeca558997dcd5d30fe21d2c256834ff Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 18:32:38 +0200 Subject: [PATCH 05/16] docs(pt): PT-2 M3 status - GGX landed, specular NEE rolls forward --- docs/pt/PT-2-real-hit-shading.md | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/docs/pt/PT-2-real-hit-shading.md b/docs/pt/PT-2-real-hit-shading.md index 7b1d889..8e1f0a2 100644 --- a/docs/pt/PT-2-real-hit-shading.md +++ b/docs/pt/PT-2-real-hit-shading.md @@ -42,10 +42,19 @@ data at every bounce: - Found + fixed the transposed `inv_vp` upload (see the post-scriptum in PT-1's ticket) — this had silently broken all PT-1 transport. -## Open (M3) +## M3 — GGX (landed b146b6d) -- GGX specular bounce using `mat_params` (port bloom-reference - `sample_brdf`: VNDF + Burley, Fresnel lobe pick). +`sample_brdf` ported: VNDF specular + Burley diffuse, Fresnel-weighted +lobe pick, metal/dielectric f0 split. Primary material from the +G-buffer material RT, bounce material from `mat_params`. Verified: the +45 s converged title render is clean (no fireflies, no NaN blowups), +with visibly deeper foliage shading than the pure-Lambert version. + +## Open (rolls into PT-3/PT-5) + +- Specular NEE (direct highlights from sun/point lights on smooth + surfaces) — bounce sampling covers reflections but not delta-light + highlights; diffuse NEE is scaled by (1 − metallic) meanwhile. - Emissive triangles at hits already work via `albedo × emissive_luma`; emissive-as-light-source sampling is PT-4 (ReSTIR) territory. - Known gaps, accepted: skinned enemies not in the TLAS (no per-frame From 90682511287cb3d08c71a3baeeffc786f19a2eac Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 19:23:15 +0200 Subject: [PATCH 06/16] fix(pt): make F9 usable in gameplay - dissolve-in progressive + firefly clamp Progressive mode (first F9) was designed for a static camera; in gameplay every camera move reset accumulation, so the screen showed raw 1-spp Monte Carlo noise through TSR - unusable. Interim behaviour until PT-3 lands the denoiser: - Mode 1 keeps the RASTER frame (with SSGI/SSR - new pt_owns_frame() gate distinct from pt_active()) while the camera moves, and skips the kernel dispatch entirely, so moving costs nothing. Hold still for ~half a second (8 samples) and the path-traced image dissolves in, refining toward ground truth. - The firefly cap now scales with accumulation depth (4 at 1 spp up to 32 converged): a lone bright GGX specular sample no longer reads as a white dot in realtime mode or during warmup. Verified: static title still converges identically; F9 cycle off->PROG->RT->off restores raster cleanly at raster fps. --- native/shared/src/renderer/mod.rs | 18 ++++++++++++++++++ native/shared/src/renderer/postfx_chain.rs | 2 +- native/shared/src/renderer/pt_pass.rs | 15 +++++++++++++++ native/shared/src/renderer/shaders/pt.rs | 17 +++++++++++++++-- native/shared/src/renderer/ssgi_pass.rs | 6 ++++-- native/shared/src/renderer/ssr_pass.rs | 4 ++-- 6 files changed, 55 insertions(+), 7 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 8caf97c..a3a3286 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -1293,6 +1293,13 @@ pub struct Renderer { /// Debug-16 numeric readback (see pt_pass.rs). pt_readback_buffer: Option, pt_dump_written: bool, + /// True when the PT kernel actually replaced this frame's scene + /// colour. Mode 1 leaves the raster frame on screen until a few + /// samples have accumulated (a moving camera = perpetual resets = + /// raw 1-spp noise), and for those frames SSGI/SSR must keep + /// running — so the downstream gates check pt_owns_frame(), not + /// pt_active(). + pt_wrote_frame: bool, /// Ticket 014 V3 — SW SDF sphere-trace pipeline + layout. /// Active whenever the scene clipmap has been baked; chosen at @@ -6790,6 +6797,7 @@ impl Renderer { pt_bg_texture_count: 0, pt_readback_buffer: None, pt_dump_written: false, + pt_wrote_frame: false, probe_trace_sdf_pipeline, probe_trace_sdf_layout, probe_trace_sdf_bg_cache: [None, None], @@ -7556,6 +7564,16 @@ impl Renderer { self.pt_mode > 0 && self.pt_supported() } + /// Whether the path tracer replaced THIS frame's scene colour. + /// Distinct from `pt_active()`: in progressive mode the kernel + /// leaves the raster frame on screen until a few samples have + /// accumulated, and those frames still need SSGI/SSR. Valid only + /// after the "pt" graph node has run for the frame (SSR/SSGI/ + /// compose all schedule after it). + fn pt_owns_frame(&self) -> bool { + self.pt_active() && self.pt_wrote_frame + } + /// Ticket 013 — rasterise every pending mesh card into its /// assigned atlas slot. Called once per frame before the HW probe /// trace samples the atlas. Drains `scene.pending_card_captures`; diff --git a/native/shared/src/renderer/postfx_chain.rs b/native/shared/src/renderer/postfx_chain.rs index 8a4845b..c6a531a 100644 --- a/native/shared/src/renderer/postfx_chain.rs +++ b/native/shared/src/renderer/postfx_chain.rs @@ -185,7 +185,7 @@ impl Renderer { // PT-1: while path tracing, the SSR passes are skipped entirely, // so history is stale — route compose to ssr_rt, which the march // else-branch keeps cleared to transparent black. - let ssr_composite_view = if self.ssr_enabled && !self.pt_active() { + let ssr_composite_view = if self.ssr_enabled && !self.pt_owns_frame() { &self.ssr_history_views[self.ssr_history_idx] } else { &self.ssr_rt_view diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index a8d9ce5..c458329 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -18,6 +18,7 @@ impl Renderer { // Leaving PT (or never entering it) invalidates history so // re-enabling starts a fresh accumulation, not a stale one. self.pt_accum_count = 0; + self.pt_wrote_frame = false; return; } // Same readiness gate as the HW probe trace: first frames before @@ -29,6 +30,7 @@ impl Renderer { || self.pt_geo_index_buffer.is_none() { self.pt_accum_count = 0; + self.pt_wrote_frame = false; return; } // PT-2 — a grown texture store means the baked view array is @@ -71,6 +73,14 @@ impl Renderer { self.pt_last_tlas_version = self.tlas_built_version; self.pt_accum_count = 0; } + // Progressive mode + camera in motion: the raster frame stays on + // screen (kernel write threshold) and any sample traced now is + // discarded by next frame's reset — skip the dispatch entirely. + // Moving costs nothing; standing still starts the accumulation. + if self.pt_mode == 1 && moved { + self.pt_wrote_frame = false; + return; + } // ---- accumulation buffer (vec4 per pixel) ---- let needed = (surf_w as u64) * (surf_h as u64) * 16; @@ -223,6 +233,11 @@ impl Renderer { } pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); } + // Mirrors the kernel's write threshold: mode 1 leaves the raster + // frame on screen until 8 samples exist (u.size.w carried the + // pre-increment count), so SSGI/SSR must keep running for those + // frames — the gates downstream check pt_owns_frame(). + self.pt_wrote_frame = self.pt_mode >= 2 || self.pt_accum_count >= 8; self.pt_accum_count = self.pt_accum_count.saturating_add(1); // ---- debug 16: numeric readback of traced intersections ---- diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index bd9e212..283d04a 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -836,12 +836,15 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } // NaN/Inf guard + firefly cap so one bad sample cannot poison the - // accumulator forever. + // accumulator forever. The cap scales with history: at 1 spp a + // single bright specular sample reads as a white dot on screen, so + // clamp hard; as the average deepens it can absorb real energy. if (radiance.r != radiance.r || radiance.g != radiance.g || radiance.b != radiance.b) { radiance = vec3(0.0); } let luma = dot(radiance, vec3(0.2126, 0.7152, 0.0722)); - if (luma > 32.0) { radiance *= 32.0 / luma; } + let cap = 4.0 + f32(min(u.size.w, 28u)); + if (luma > cap) { radiance *= cap / luma; } // ---- accumulate --------------------------------------------------------- @@ -863,6 +866,16 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let n = f32(u.size.w) + 1.0; accum[idx] = vec4(sum, n); out = sum / n; + // Interim gameplay behaviour until PT-3's denoiser: a moving + // camera resets accumulation every frame, and raw 1-spp noise + // through TSR looks terrible. Keep accumulating but leave the + // raster frame on screen until a few samples exist — stand + // still for half a second and PT dissolves in. The CPU side + // mirrors this threshold (pt_wrote_frame) so SSGI/SSR stay on + // for the raster frames. + if (u.size.w < 8u) { + return; + } } textureStore(out_hdr, px, vec4(out, 1.0)); diff --git a/native/shared/src/renderer/ssgi_pass.rs b/native/shared/src/renderer/ssgi_pass.rs index f4f5e54..62a2c40 100644 --- a/native/shared/src/renderer/ssgi_pass.rs +++ b/native/shared/src/renderer/ssgi_pass.rs @@ -30,8 +30,10 @@ impl Renderer { // PT-1: while the path tracer owns the frame its output already // contains full GI — probe SSGI would burn ~2ms to be composited // over by nothing (the else-branch clear keeps compose additive- - // safe either way). - if self.ssgi_enabled && !self.pt_active() { + // safe either way). pt_owns_frame, not pt_active: progressive mode + // shows raster frames until accumulation warms up, and those still + // want SSGI. + if self.ssgi_enabled && !self.pt_owns_frame() { let p00 = self.current_proj_matrix[0][0]; let p11 = self.current_proj_matrix[1][1]; let p20 = self.current_proj_matrix[2][0]; diff --git a/native/shared/src/renderer/ssr_pass.rs b/native/shared/src/renderer/ssr_pass.rs index a56f992..5a09a8c 100644 --- a/native/shared/src/renderer/ssr_pass.rs +++ b/native/shared/src/renderer/ssr_pass.rs @@ -17,7 +17,7 @@ impl Renderer { // ============================================================ // PT-1: skipped while the path tracer owns the frame — it marches // the raster-lit HDR, which PT has already overwritten. - if self.ssr_enabled && !self.pt_active() { + if self.ssr_enabled && !self.pt_owns_frame() { let inv_proj = self.current_inv_proj_matrix; // EN-021 — view→world rotation for the env-miss fallback: the // transpose of the view matrix's 3×3 (rigid view ⇒ inverse @@ -139,7 +139,7 @@ impl Renderer { // ssr_rt. // ============================================================ // PT-1: same gate as the march — no fresh rays, nothing to blend. - if self.ssr_enabled && !self.pt_active() { + if self.ssr_enabled && !self.pt_owns_frame() { let prev_idx = 1 - self.ssr_history_idx; let cur_idx = self.ssr_history_idx; From 80be92a3e6edd2b2da56328de01979798e69de3f Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 20:23:18 +0200 Subject: [PATCH 07/16] feat(pt): PT-3 M1 - temporally reprojected realtime mode + AO degate RT mode (F9 x2) blended the last ~8 frames per SCREEN PIXEL with no reprojection, so any camera motion averaged different viewpoints into smear - the foggy-lens look. History now follows world points: the previous frame result is fetched through prev_vp (unjittered, transposed like inv_vp) into a ping-pong accumulation pair, validated by depth (accum.w carries the raw depth its surface wrote; mismatch or off-screen rejects history, alpha=1). Verified crisp mid-strafe via a held-PostMessage-key harness - held keys DO reach game input. Also degates screen-space AO while PT owns the frame (ssao_blur white-clears): PT computes real occlusion by tracing, and GTAO on top double-darkened every crevice - one reason converged PT looked barely different from raster. Remaining PT-3: a-trous spatial filter over the residual grain + half-res tracing for the fps budget. --- native/shared/src/renderer/hiz.rs | 6 ++- native/shared/src/renderer/mod.rs | 48 ++++++++++++----- native/shared/src/renderer/pt_pass.rs | 66 +++++++++++++++++------- native/shared/src/renderer/shaders/pt.rs | 55 +++++++++++++++++--- 4 files changed, 132 insertions(+), 43 deletions(-) diff --git a/native/shared/src/renderer/hiz.rs b/native/shared/src/renderer/hiz.rs index 5db7447..e5fa3a0 100644 --- a/native/shared/src/renderer/hiz.rs +++ b/native/shared/src/renderer/hiz.rs @@ -104,7 +104,11 @@ impl Renderer { // preserving depth edges (depth-guided bilateral filter). // Reads ssao_rt → writes ssao_blur_rt. // ============================================================ - if self.ssao_enabled { + // PT: when the path tracer owns the frame it computes real + // occlusion by tracing — screen-space AO on top double-darkens + // every crevice. Route compose to "no occlusion" (white clear + // below) for those frames. + if self.ssao_enabled && !self.pt_owns_frame() { // texel_size is the size of one SSAO RT texel (half-res). let ao_w = (surf_w / 2).max(1) as f32; let ao_h = (surf_h / 2).max(1) as f32; diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index a3a3286..a1c2a83 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -370,6 +370,9 @@ const PT_MAX_TEXTURES: usize = 256; #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] struct PtParamsCpu { inv_vp: [[f32; 4]; 4], + /// PT-3 — previous frame's unjittered VP (transposed like inv_vp) + /// for temporal history reprojection in realtime mode. + prev_vp: [[f32; 4]; 4], /// xyz camera world pos, w unused. cam_pos: [f32; 4], /// xyz unit vector toward the sun, w unused. @@ -1259,12 +1262,19 @@ pub struct Renderer { pub pt_pipeline: Option, pub pt_layout: Option, /// Rebuilt lazily; nulled on resize and on instance-data rebuild. - pt_bg: Option, + /// Two variants for the accumulation ping-pong: bg[i] reads + /// accum_buffers[i] (binding 8) and writes accum_buffers[1-i] + /// (binding 13). + pt_bg: [Option; 2], pt_uniform_buffer: wgpu::Buffer, - /// rgba32float-equivalent accumulation (vec4 per pixel) as a storage - /// buffer — rgba32float storage *textures* lack read_write in core - /// WGSL. Lazily (re)created at render extent. - pt_accum_buffer: Option, + /// rgba32float-equivalent accumulation (vec4 per pixel) as storage + /// buffers — rgba32float storage *textures* lack read_write in core + /// WGSL. Ping-pong pair (PT-3 reprojection reads other pixels from + /// the previous frame). Lazily (re)created at render extent. + pt_accum_buffers: [Option; 2], + /// Index of the buffer holding the PREVIOUS frame's result (the + /// read side of this frame's dispatch). Flipped after dispatch. + pt_accum_idx: usize, /// Samples accumulated at the current view; 0 = history invalid. pt_accum_count: u32, /// VP matrix of the last accumulated frame; movement beyond epsilon @@ -5126,6 +5136,15 @@ impl Renderer { has_dynamic_offset: false, min_binding_size: None, }, count: None, }, + // PT-3 — accumulation write target (binding 8 holds + // the previous frame's buffer; ping-pong). + wgpu::BindGroupLayoutEntry { + binding: 13, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, ]; let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("pt_layout"), @@ -6782,9 +6801,10 @@ impl Renderer { probe_trace_hw_bg_cache: [None, None], pt_pipeline, pt_layout, - pt_bg: None, + pt_bg: [None, None], pt_uniform_buffer, - pt_accum_buffer: None, + pt_accum_buffers: [None, None], + pt_accum_idx: 0, pt_accum_count: 0, pt_prev_vp: [[0.0; 4]; 4], pt_last_tlas_version: 0, @@ -7293,11 +7313,11 @@ impl Renderer { self.probe_trace_sdf_bg_cache = [None, None]; self.probe_temporal_bg_cache = [None, None]; self.probe_resolve_bg_cache = [None, None]; - // PT-1 — the BG references hdr_rt/depth views and the - // accumulation buffer is per-pixel-sized; both die with + // PT-1 — the BGs reference hdr_rt/depth views and the + // accumulation buffers are per-pixel-sized; both die with // the old extent. - self.pt_bg = None; - self.pt_accum_buffer = None; + self.pt_bg = [None, None]; + self.pt_accum_buffers = [None, None]; self.pt_accum_count = 0; self.hiz_linearize_bg_cache = None; self.occlusion.invalidate_bindings(); @@ -7995,7 +8015,7 @@ impl Renderer { // instance_data buffer; invalidate on resize. self.wsrc_bake_hw_bg_cache = None; // PT-1 — same buffer bound at group 0 binding 2. - self.pt_bg = None; + self.pt_bg = [None, None]; resized = true; } @@ -8091,7 +8111,7 @@ impl Renderer { })); } if v_recreate || i_recreate { - self.pt_bg = None; + self.pt_bg = [None, None]; } if !geo_vertices.is_empty() { self.queue.write_buffer( @@ -8188,7 +8208,7 @@ impl Renderer { // V14 — same reason as the resize path above. self.wsrc_bake_hw_bg_cache = None; // PT-1 — TLAS bound at group 0 binding 1. - self.pt_bg = None; + self.pt_bg = [None, None]; } // Populate TLAS instance slots. Clear stale entries from prior diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index c458329..0b23807 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -66,6 +66,9 @@ impl Renderer { if moved && self.pt_mode == 1 { self.pt_accum_count = 0; } + // PT-3 — the uniform needs LAST frame's VP for history + // reprojection; stash it before the tracker is overwritten. + let prev_vp_for_reproject = self.pt_prev_vp; self.pt_prev_vp = vp_unjittered; // Geometry changed under the accumulated image (door opened, // enemy died) → history is a lie, restart. @@ -82,25 +85,28 @@ impl Renderer { return; } - // ---- accumulation buffer (vec4 per pixel) ---- + // ---- accumulation buffers (vec4 per pixel, ping-pong) ---- let needed = (surf_w as u64) * (surf_h as u64) * 16; - let recreate = match &self.pt_accum_buffer { + let recreate = match &self.pt_accum_buffers[0] { Some(b) => b.size() != needed, None => true, }; if recreate { - self.pt_accum_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { - label: Some("pt_accum"), - size: needed, - // COPY_SRC: the debug-16 numeric readback copies a window - // of this buffer to a staging buffer. - usage: wgpu::BufferUsages::STORAGE - | wgpu::BufferUsages::COPY_DST - | wgpu::BufferUsages::COPY_SRC, - mapped_at_creation: false, - })); - self.pt_bg = None; + for (i, slot) in self.pt_accum_buffers.iter_mut().enumerate() { + *slot = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some(if i == 0 { "pt_accum_a" } else { "pt_accum_b" }), + size: needed, + // COPY_SRC: the debug-16 numeric readback copies a + // window of this buffer to a staging buffer. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + })); + } + self.pt_bg = [None, None]; self.pt_accum_count = 0; + self.pt_accum_idx = 0; } // ---- uniforms ---- @@ -137,8 +143,17 @@ impl Renderer { [m[0][2], m[1][2], m[2][2], m[3][2]], [m[0][3], m[1][3], m[2][3], m[3][3]], ]; + // Same transpose story for the reprojection matrix. + let p = &prev_vp_for_reproject; + let prev_vp_t = [ + [p[0][0], p[1][0], p[2][0], p[3][0]], + [p[0][1], p[1][1], p[2][1], p[3][1]], + [p[0][2], p[1][2], p[2][2], p[3][2]], + [p[0][3], p[1][3], p[2][3], p[3][3]], + ]; let params = PtParamsCpu { inv_vp: inv_vp_t, + prev_vp: prev_vp_t, cam_pos: [cam[0], cam[1], cam[2], 0.0], sun_dir: [ -ld[0] * sun_inv_len, @@ -169,9 +184,13 @@ impl Renderer { }; self.queue.write_buffer(&self.pt_uniform_buffer, 0, bytemuck::bytes_of(¶ms)); - // ---- bind group (lazy; nulled on resize / TLAS or instance - // buffer recreation) ---- - if self.pt_bg.is_none() { + // ---- bind groups (lazy; nulled on resize / TLAS or instance + // buffer recreation). Two ping-pong variants: bg[i] reads accum + // buffer i (binding 8) and writes buffer 1-i (binding 13). + for i in 0..2 { + if self.pt_bg[i].is_some() { + continue; + } let tlas = self.tlas.as_ref().unwrap(); let entries = vec![ wgpu::BindGroupEntry { binding: 0, resource: self.pt_uniform_buffer.as_entire_binding() }, @@ -185,12 +204,13 @@ impl Renderer { // hits; radiance would double-count. wgpu::BindGroupEntry { binding: 6, resource: wgpu::BindingResource::TextureView(&self.mesh_card_atlas_view) }, wgpu::BindGroupEntry { binding: 7, resource: wgpu::BindingResource::Sampler(&self.mesh_card_atlas_sampler) }, - wgpu::BindGroupEntry { binding: 8, resource: self.pt_accum_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 8, resource: self.pt_accum_buffers[i].as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 9, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, wgpu::BindGroupEntry { binding: 10, resource: self.pt_geo_vertex_buffer.as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 11, resource: self.pt_geo_index_buffer.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 13, resource: self.pt_accum_buffers[1 - i].as_ref().unwrap().as_entire_binding() }, ]; - self.pt_bg = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + self.pt_bg[i] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("pt_bg"), layout: self.pt_layout.as_ref().unwrap(), entries: &entries, @@ -227,12 +247,16 @@ impl Renderer { timestamp_writes: ts, }); pass.set_pipeline(self.pt_pipeline.as_ref().unwrap()); - pass.set_bind_group(0, self.pt_bg.as_ref().unwrap(), &[]); + pass.set_bind_group(0, self.pt_bg[self.pt_accum_idx].as_ref().unwrap(), &[]); if self.pt_texture_arrays_enabled { pass.set_bind_group(1, self.pt_tex_bg.as_ref().unwrap(), &[]); } pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); } + // This frame wrote into buffers[1 - idx]; it becomes next + // frame's read side. + let written_idx = 1 - self.pt_accum_idx; + self.pt_accum_idx = written_idx; // Mirrors the kernel's write threshold: mode 1 leaves the raster // frame on screen until 8 samples exist (u.size.w carried the // pre-increment count), so SSGI/SSR must keep running for those @@ -259,7 +283,9 @@ impl Renderer { mapped_at_creation: false, })); encoder.copy_buffer_to_buffer( - self.pt_accum_buffer.as_ref().unwrap(), + // written_idx == pt_accum_idx here (already flipped): + // the buffer this frame's dispatch wrote. + self.pt_accum_buffers[self.pt_accum_idx].as_ref().unwrap(), offset, self.pt_readback_buffer.as_ref().unwrap(), 0, diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 283d04a..d2f7b3b 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -47,6 +47,10 @@ struct PtLight { struct PtParams { inv_vp: mat4x4, + // PT-3: previous frame's UNJITTERED view-projection — reprojects + // this frame's world positions into last frame's screen for + // temporal history fetch in realtime mode. + prev_vp: mat4x4, cam_pos: vec4, // xyz camera world pos sun_dir: vec4, // xyz unit vector toward the sun sun_color: vec4, // rgb premultiplied by intensity @@ -82,8 +86,19 @@ struct InstanceGiData { @group(0) @binding(5) var material_tex: texture_2d; @group(0) @binding(6) var card_albedo_atlas: texture_2d; @group(0) @binding(7) var card_samp: sampler; +// PT-3: ping-pong accumulation. Binding 8 = previous frame's buffer +// (read), binding 13 = this frame's output. Reprojection reads OTHER +// pixels from prev, which a single read_write buffer cannot do safely. @group(0) @binding(8) var accum: array>; @group(0) @binding(9) var out_hdr: texture_storage_2d; +@group(0) @binding(13) var accum_out: array>; + +// Approximate linear view distance from the raw depth-buffer value +// (GL-convention matrix, near 0.01: z_view ~= 2n / (1 - d)). Only used +// for RELATIVE history-validation comparisons. +fn lin_depth(d: f32) -> f32 { + return 0.02 / max(1.0 - d, 1e-6); +} // PT-2: geometry megabuffers. geo_v holds raw Vertex3D words (stride 24 // f32: position +0, normal +3, color +6, uv +10, ...); geo_i holds the // concatenated index streams. Windows are per-instance via inst.geo. @@ -698,7 +713,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } let h16 = rayQueryGetCommittedIntersection(&rq16); let idx16 = gid.y * u.size.x + gid.x; - accum[idx16] = vec4( + accum_out[idx16] = vec4( h16.t, f32(h16.instance_custom_data), f32(h16.primitive_index), @@ -711,7 +726,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // Raw ray-generation dump: reconstructed world position + raw // depth, straight into accum for CPU readback. let idx17 = gid.y * u.size.x + gid.x; - accum[idx17] = vec4(p0, depth); + accum_out[idx17] = vec4(p0, depth); textureStore(out_hdr, px, vec4(0.4, 0.2, 0.0, 1.0)); return; } @@ -855,16 +870,40 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { var out: vec3; if (mode >= 2.0) { - // Realtime: rolling exponential window. PT-3 replaces this with - // motion-reprojected accumulation + a-trous. - let alpha = select(0.125, 1.0, u.size.w == 0u); - out = mix(prev.rgb, radiance, alpha); - accum[idx] = vec4(out, 1.0); + // PT-3 M1: temporally REPROJECTED exponential history. The + // previous frame's result is fetched where THIS surface was + // last frame (world pos through prev_vp), so camera motion no + // longer smears screen-space history into fog. Disocclusions + // and depth mismatches reject history (alpha = 1); accepted + // history blends at 0.15. accum.w carries the raw depth the + // texel's surface wrote, validated against the reprojected + // depth of this frame's surface. + var alpha = 1.0; + var hist = vec3(0.0); + let clip_prev = u.prev_vp * vec4(p0, 1.0); + if (u.size.w > 0u && clip_prev.w > 1e-4) { + let ndc_prev = clip_prev.xyz / clip_prev.w; + let uv_prev = vec2(ndc_prev.x * 0.5 + 0.5, 0.5 - ndc_prev.y * 0.5); + if (uv_prev.x >= 0.0 && uv_prev.x < 1.0 && uv_prev.y >= 0.0 && uv_prev.y < 1.0) { + let ppx = vec2(uv_prev * vec2(f32(u.size.x), f32(u.size.y))); + let pidx = min(ppx.y, u.size.y - 1u) * u.size.x + min(ppx.x, u.size.x - 1u); + let ph = accum[pidx]; + let zl_hist = lin_depth(ph.w); + let zl_here = lin_depth(ndc_prev.z); + if (abs(zl_hist - zl_here) < 0.05 * max(zl_hist, zl_here) + 0.05) { + alpha = 0.15; + hist = ph.rgb; + } + } + } + out = mix(hist, radiance, alpha); + accum_out[idx] = vec4(out, depth); } else { // Progressive: plain running sum; count lives on the CPU. + // Ping-pong read/write at the same index (static camera only). let sum = prev.rgb + radiance; let n = f32(u.size.w) + 1.0; - accum[idx] = vec4(sum, n); + accum_out[idx] = vec4(sum, n); out = sum / n; // Interim gameplay behaviour until PT-3's denoiser: a moving // camera resets accumulation every frame, and raw 1-spp noise From 8f239295f6cd82904da4f2a8edc383b0c072c450 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 23:26:20 +0200 Subject: [PATCH 08/16] fix(pt): kill RT-mode fireflies at the source - Fresnel-at-view lobe pick The GGX lobe was picked with the Fresnel-at-normal weight (~4% on dielectrics) and the estimator divides by that probability - but real Fresnel approaches 1 at grazing angles, so rare grazing specular picks came back amplified ~25x. The whole ground plane is grazing at distance: at 1-2 spp that reads as white salt across the frame. Picking by Fresnel at the actual view angle makes the divisor track the true specular energy (clamped [0.05, 0.95] so neither lobe boosts past ~20x). Realtime mode additionally clamps per-bounce weights to 4x and uses a FIXED radiance cap of 6 - its EMA window stays ~7 frames no matter how large the frame counter grows, so the count-scaled cap had quietly relaxed back to 32 and stopped protecting anything. Progressive mode keeps unclamped weights + the count-scaled cap (deep averages absorb real energy; bias not needed there). Verified mid-strafe: uniform low-contrast grain, no white dots. --- native/shared/src/renderer/shaders/pt.rs | 35 ++++++++++++++++++++---- 1 file changed, 30 insertions(+), 5 deletions(-) diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index d2f7b3b..bce6a08 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -333,9 +333,19 @@ fn sample_brdf( return out; } let f0 = mix(vec3(0.04), base_color, metallic); - let spec_weight = (f0.x + f0.y + f0.z) / 3.0; + // Lobe pick by Fresnel at the ACTUAL view angle, not at normal + // incidence: at grazing angles specular energy approaches 1, and + // the estimator divides by the pick probability — picking with the + // ~0.04 normal-incidence weight amplified rare grazing specular + // samples ~25x into a field of white fireflies at 1-2 spp (the + // whole ground plane is grazing at distance). Clamped so neither + // lobe's 1/p boost can exceed ~20x even in edge cases. + let n_dot_v_pick = max(dot(n, view_ws), 0.0); + let f_view = fresnel_schlick3(n_dot_v_pick, f0); + let spec_weight = (f_view.x + f_view.y + f_view.z) / 3.0; let diff_weight = (1.0 - spec_weight) * (1.0 - metallic); - let p_spec = spec_weight / (spec_weight + diff_weight + 1e-6); + var p_spec = spec_weight / (spec_weight + diff_weight + 1e-6); + p_spec = clamp(p_spec, 0.05, 0.95); let r2 = rand_2f(); if (rand_f() < p_spec) { let h_t = sample_ggx_vndf(v_t, alpha, r2); @@ -351,7 +361,14 @@ fn sample_brdf( let g2 = v_smith(n_dot_v, n_dot_l, alpha) * 4.0 * n_dot_v * n_dot_l; let g1_v = smith_g1(n_dot_v, alpha); out.dir = m * l_t; - out.weight = f * g2 / (max(g1_v, 1e-6) * max(p_spec, 1e-6)); + out.weight = f * g2 / (max(g1_v, 1e-6) * p_spec); + // Realtime mode trades a little energy for stability: a single + // bounce may not multiply throughput more than 4x (the ~7-frame + // EMA window cannot average outliers away like progressive + // accumulation can). Progressive mode stays unclamped. + if (u.cfg.x >= 2.0) { + out.weight = min(out.weight, vec3(4.0)); + } out.valid = true; return out; } @@ -370,7 +387,10 @@ fn sample_brdf( let diffuse_albedo = base_color * (1.0 - metallic) * (vec3(1.0) - f0); let fd = burley_diffuse(n_dot_l, n_dot_v, l_dot_h, roughness); out.dir = m * l_t; - out.weight = diffuse_albedo * fd * 3.14159265 / max(1.0 - p_spec, 1e-6); + out.weight = diffuse_albedo * fd * 3.14159265 / (1.0 - p_spec); + if (u.cfg.x >= 2.0) { + out.weight = min(out.weight, vec3(4.0)); + } out.valid = true; return out; } @@ -858,7 +878,12 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { radiance = vec3(0.0); } let luma = dot(radiance, vec3(0.2126, 0.7152, 0.0722)); - let cap = 4.0 + f32(min(u.size.w, 28u)); + // Progressive: relax with accumulation depth (a deep average can + // absorb real energy). Realtime: FIXED — its EMA window stays ~7 + // frames no matter how large the frame count grows, so the cap + // must never relax or fireflies return as the count climbs. + var cap = 4.0 + f32(min(u.size.w, 28u)); + if (u.cfg.x >= 2.0) { cap = 6.0; } if (luma > cap) { radiance *= cap / luma; } // ---- accumulate --------------------------------------------------------- From a0b91156800fa081b9eb469e18d02284d209120e Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Mon, 13 Jul 2026 23:42:48 +0200 Subject: [PATCH 09/16] feat(pt): PT-3 M2 - edge-aware a-trous denoiser for realtime mode Two 5x5 B3-spline a-trous iterations (step 1, then 2) over the temporally-blended radiance, edge-stopped by relative linearized depth (accum.w) and luminance. The main kernel now only feeds the history buffer in realtime mode; cs_mid filters accum->scratch and cs_final filters scratch->hdr, skipping sky texels (w=1 marker) so the raster sky survives. Sigma_luma 0.25, both dispatches in one compute pass. This addresses the user-visible boiling: per-frame 1-2 spp grain re-rolls every frame and reads as crawling white shimmer live even when stills look acceptable. Filtered mid-strafe capture is smooth with a faint dusting - the remaining lever is half-res tracing for frame rate (RT ~11 fps at 1080p internal on the 760M). --- native/shared/src/renderer/mod.rs | 102 +++++++++++++++++++++- native/shared/src/renderer/pt_pass.rs | 65 ++++++++++++++ native/shared/src/renderer/shaders/mod.rs | 2 +- native/shared/src/renderer/shaders/pt.rs | 101 ++++++++++++++++++++- 4 files changed, 267 insertions(+), 3 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index a1c2a83..e8ff774 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -1303,6 +1303,18 @@ pub struct Renderer { /// Debug-16 numeric readback (see pt_pass.rs). pt_readback_buffer: Option, pt_dump_written: bool, + /// PT-3 — à-trous denoiser (realtime mode): two iterations over the + /// temporally-blended radiance, buffer→scratch then scratch→hdr. + pt_atrous_mid_pipeline: Option, + pt_atrous_final_pipeline: Option, + pt_atrous_layout: Option, + /// Per-iteration uniform (step size + sigmas + extent). + pt_atrous_params_bufs: [wgpu::Buffer; 2], + pt_atrous_scratch: Option, + /// bg[0][i] = pass 1 reading accum buffer i; bg[1][0] = pass 2 + /// (scratch → hdr). Invalidated with pt_bg. + pt_atrous_bg1: [Option; 2], + pt_atrous_bg2: Option, /// True when the PT kernel actually replaced this frame's scene /// colour. Mode 1 leaves the raster frame on screen until a few /// samples have accumulated (a moving camera = perpetual resets = @@ -5189,6 +5201,84 @@ impl Renderer { } else { (None, None, None) }; + // PT-3 — à-trous denoiser pipelines for the realtime mode + // (plain compute; gated with the kernel since it is useless + // without it). + let (pt_atrous_mid_pipeline, pt_atrous_final_pipeline, pt_atrous_layout) = + if hw_rt_enabled { + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("pt_atrous_shader"), + source: wgpu::ShaderSource::Wgsl(PT_ATROUS_WGSL.into()), + }); + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("pt_atrous_layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 3, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: HDR_FORMAT, + view_dimension: wgpu::TextureViewDimension::D2, + }, count: None, + }, + ], + }); + let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("pt_atrous_pl"), + bind_group_layouts: &[Some(&layout)], + immediate_size: 0, + }); + let mid = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("pt_atrous_mid"), + layout: Some(&pl), + module: &shader, entry_point: Some("cs_mid"), + compilation_options: Default::default(), cache: None, + }); + let fin = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("pt_atrous_final"), + layout: Some(&pl), + module: &shader, entry_point: Some("cs_final"), + compilation_options: Default::default(), cache: None, + }); + (Some(mid), Some(fin), Some(layout)) + } else { + (None, None, None) + }; + let pt_atrous_params_bufs = [ + device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_atrous_params_0"), + size: 16, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_atrous_params_1"), + size: 16, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + ]; let pt_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_uniform"), size: std::mem::size_of::() as u64, @@ -6818,6 +6908,13 @@ impl Renderer { pt_readback_buffer: None, pt_dump_written: false, pt_wrote_frame: false, + pt_atrous_mid_pipeline, + pt_atrous_final_pipeline, + pt_atrous_layout, + pt_atrous_params_bufs, + pt_atrous_scratch: None, + pt_atrous_bg1: [None, None], + pt_atrous_bg2: None, probe_trace_sdf_pipeline, probe_trace_sdf_layout, probe_trace_sdf_bg_cache: [None, None], @@ -7315,10 +7412,13 @@ impl Renderer { self.probe_resolve_bg_cache = [None, None]; // PT-1 — the BGs reference hdr_rt/depth views and the // accumulation buffers are per-pixel-sized; both die with - // the old extent. + // the old extent. Same for the à-trous scratch + bgs. self.pt_bg = [None, None]; self.pt_accum_buffers = [None, None]; self.pt_accum_count = 0; + self.pt_atrous_scratch = None; + self.pt_atrous_bg1 = [None, None]; + self.pt_atrous_bg2 = None; self.hiz_linearize_bg_cache = None; self.occlusion.invalidate_bindings(); for slot in self.hiz_downsample_bg_cache.iter_mut() { diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index 0b23807..b92bb6e 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -104,7 +104,15 @@ impl Renderer { mapped_at_creation: false, })); } + self.pt_atrous_scratch = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_atrous_scratch"), + size: needed, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + })); self.pt_bg = [None, None]; + self.pt_atrous_bg1 = [None, None]; + self.pt_atrous_bg2 = None; self.pt_accum_count = 0; self.pt_accum_idx = 0; } @@ -257,6 +265,63 @@ impl Renderer { // frame's read side. let written_idx = 1 - self.pt_accum_idx; self.pt_accum_idx = written_idx; + + // ---- PT-3: à-trous denoise (realtime mode only) ---- + // Two edge-aware iterations over the temporally-blended + // radiance: accum[written] → scratch (step 1), scratch → hdr + // (step 2). Progressive mode converges on its own and writes + // hdr directly from the kernel. + if self.pt_mode >= 2 + && self.pt_atrous_mid_pipeline.is_some() + && self.pt_atrous_scratch.is_some() + { + let sigma_luma = 0.25f32; + let p0 = [1.0f32, sigma_luma, surf_w as f32, surf_h as f32]; + let p1 = [2.0f32, sigma_luma, surf_w as f32, surf_h as f32]; + self.queue.write_buffer(&self.pt_atrous_params_bufs[0], 0, bytemuck::bytes_of(&p0)); + self.queue.write_buffer(&self.pt_atrous_params_bufs[1], 0, bytemuck::bytes_of(&p1)); + + if self.pt_atrous_bg1[written_idx].is_none() { + self.pt_atrous_bg1[written_idx] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_atrous_bg1"), + layout: self.pt_atrous_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[0].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: self.pt_accum_buffers[written_idx].as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + ], + })); + } + if self.pt_atrous_bg2.is_none() { + self.pt_atrous_bg2 = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_atrous_bg2"), + layout: self.pt_atrous_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[1].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, + // dst is unused by cs_final; bind an accum + // buffer to satisfy the layout (cannot alias + // the scratch src binding — RO+RW of the same + // buffer in one group fails validation). + wgpu::BindGroupEntry { binding: 2, resource: self.pt_accum_buffers[0].as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + ], + })); + } + + let ts = profiler.compute_pass_timestamp_writes("pt_atrous"); + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("pt_atrous"), + timestamp_writes: ts, + }); + pass.set_pipeline(self.pt_atrous_mid_pipeline.as_ref().unwrap()); + pass.set_bind_group(0, self.pt_atrous_bg1[written_idx].as_ref().unwrap(), &[]); + pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); + pass.set_pipeline(self.pt_atrous_final_pipeline.as_ref().unwrap()); + pass.set_bind_group(0, self.pt_atrous_bg2.as_ref().unwrap(), &[]); + pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); + } // Mirrors the kernel's write threshold: mode 1 leaves the raster // frame on screen until 8 samples exist (u.size.w carried the // pre-increment count), so SSGI/SSR must keep running for those diff --git a/native/shared/src/renderer/shaders/mod.rs b/native/shared/src/renderer/shaders/mod.rs index 312e766..6e3d8d0 100644 --- a/native/shared/src/renderer/shaders/mod.rs +++ b/native/shared/src/renderer/shaders/mod.rs @@ -14,6 +14,6 @@ pub(super) use gi::{CARD_CAPTURE_WGSL, SDF_BAKE_WGSL, SDF_CLIPMAP_BAKE_WGSL, CAR mod ssgi; pub(super) use ssgi::{PROBE_HELPERS_WGSL, SSGI_PROBE_PLACE_WGSL, SSGI_PROBE_TRACE_SW_WGSL, SSGI_PROBE_TRACE_HW_WGSL, SSGI_PROBE_TRACE_SDF_WGSL, SSGI_PROBE_TEMPORAL_WGSL, SSGI_PROBE_RESOLVE_WGSL, SSR_TEMPORAL_SHADER_WGSL, SSR_SHADER_WGSL}; mod pt; -pub(super) use pt::PT_KERNEL_WGSL; +pub(super) use pt::{PT_KERNEL_WGSL, PT_ATROUS_WGSL}; mod post; pub(super) use post::{BLOOM_SHADER_WGSL, DOF_SHADER_WGSL, MOTION_BLUR_SHADER_WGSL, SSS_SHADER_WGSL, SCENE_COMPOSE_SHADER_WGSL, TAA_SHADER_WGSL, EXPOSURE_SHADER_WGSL, COMPOSITE_SHADER_WGSL, UPSCALE_SHADER_WGSL, RCAS_SHADER_WGSL}; diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index bce6a08..1eb31e6 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -523,7 +523,11 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } if (is_sky(depth)) { - // Leave the raster sky/clouds untouched. + // Leave the raster sky/clouds untouched. Realtime mode marks + // the texel (w = 1) so the a-trous passes skip it too. + if (u.cfg.x >= 2.0 && u.cfg.w == 0.0) { + accum_out[gid.y * u.size.x + gid.x] = vec4(0.0, 0.0, 0.0, 1.0); + } return; } @@ -923,6 +927,9 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } out = mix(hist, radiance, alpha); accum_out[idx] = vec4(out, depth); + // Realtime output goes through the a-trous passes (they write + // out_hdr); the main kernel only feeds the history buffer. + return; } else { // Progressive: plain running sum; count lives on the CPU. // Ping-pong read/write at the same index (static camera only). @@ -945,3 +952,95 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { textureStore(out_hdr, px, vec4(out, 1.0)); } "#; + +/// PT-3 — edge-aware à-trous denoiser for the realtime mode. Two +/// iterations (step 1 then 2) over the temporally-blended radiance: +/// `cs_mid` filters buffer→buffer, `cs_final` filters buffer→out_hdr. +/// Edge stopping = relative linearized depth (accum.w carries the raw +/// depth) + luminance similarity; sky texels (w = 1 marker) pass +/// through untouched and are never written to hdr. +pub(in crate::renderer) const PT_ATROUS_WGSL: &str = r#" +struct AtrousParams { + // x = step (texels), y = sigma_luma, z = width, w = height + p: vec4, +}; +@group(0) @binding(0) var ap: AtrousParams; +@group(0) @binding(1) var src: array>; +@group(0) @binding(2) var dst: array>; +@group(0) @binding(3) var out_hdr_a: texture_storage_2d; + +fn lin_depth_a(d: f32) -> f32 { + return 0.02 / max(1.0 - d, 1e-6); +} + +// B3-spline kernel weight for |offset| 0/1/2. +fn kern(d: i32) -> f32 { + let a = abs(d); + if (a == 0) { return 0.375; } + if (a == 1) { return 0.25; } + return 0.0625; +} + +fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { + let cidx = u32(px.y) * u32(w) + u32(px.x); + let center = src[cidx]; + if (center.w >= 0.9999999) { + return center; + } + let zc = lin_depth_a(center.w); + let lc = dot(center.rgb, vec3(0.2126, 0.7152, 0.0722)); + var sum = vec3(0.0); + var wsum = 0.0; + for (var dy = -2; dy <= 2; dy = dy + 1) { + for (var dx = -2; dx <= 2; dx = dx + 1) { + let q = px + vec2(dx, dy) * step; + if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { + continue; + } + let s = src[u32(q.y) * u32(w) + u32(q.x)]; + if (s.w >= 0.9999999) { + continue; + } + let zq = lin_depth_a(s.w); + let wz = exp(-abs(zq - zc) / (0.08 * max(zc, zq) + 0.02)); + let lq = dot(s.rgb, vec3(0.2126, 0.7152, 0.0722)); + let wl = exp(-abs(lq - lc) / max(ap.p.y, 1e-3)); + let wgt = kern(dx) * kern(dy) * wz * wl; + sum += s.rgb * wgt; + wsum += wgt; + } + } + if (wsum < 1e-6) { + return center; + } + return vec4(sum / wsum, center.w); +} + +@compute @workgroup_size(8, 8, 1) +fn cs_mid(@builtin(global_invocation_id) gid: vec3) { + let w = i32(ap.p.z); + let h = i32(ap.p.w); + if (i32(gid.x) >= w || i32(gid.y) >= h) { + return; + } + let px = vec2(i32(gid.x), i32(gid.y)); + dst[gid.y * u32(w) + gid.x] = filter_at(px, w, h, i32(ap.p.x)); +} + +@compute @workgroup_size(8, 8, 1) +fn cs_final(@builtin(global_invocation_id) gid: vec3) { + let w = i32(ap.p.z); + let h = i32(ap.p.w); + if (i32(gid.x) >= w || i32(gid.y) >= h) { + return; + } + let px = vec2(i32(gid.x), i32(gid.y)); + let cidx = gid.y * u32(w) + gid.x; + if (src[cidx].w >= 0.9999999) { + // Sky: the raster sky in hdr stays untouched. + return; + } + let r = filter_at(px, w, h, i32(ap.p.x)); + textureStore(out_hdr_a, px, vec4(r.rgb, 1.0)); +} +"#; From 5f3ae83cd76b2b211f5fd23f0d049137ec11703c Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 00:17:20 +0200 Subject: [PATCH 10/16] feat(pt): PT-3 M3 - half-res tracing + SVGF-style filtering for RT mode Realtime mode now traces on a half-resolution grid (4x fewer rays; ~11 -> ~20+ fps mid-motion at 1080p internal on the 760M) and reconstructs full resolution in the final a-trous pass via depth-guided joint-bilateral upsampling. The pieces that made it look GOOD and not just fast, each verified with mid-strafe captures: - Albedo demodulation: history stores IRRADIANCE (radiance divided by primary albedo); the final pass re-multiplies by the full-res G-buffer albedo. Texture detail stays crisp despite the smoothing. - The firefly cap moved to irradiance space (cap 4): capping radiance before the albedo divide let dark-albedo texels amplify outliers ~20x past every filter sigma - white 2x2 blocks. - Three a-trous iterations (steps 1, 2, 4) with an SVGF-style sigma schedule: the first iteration barely edge-stops on luminance (1-spp variance IS the noise), later ones tighten to preserve shading. - Loosened history-validation + filter depth tolerances: at half-res the reprojected texel quantizes to 2 full pixels, and tight tests rejected nearly every frame on depth-chaotic surfaces (grass), pinning them at 1 spp. - Sample phase pinned to (0,0): rotating it re-rolled the owner pixel every frame and broke history validation the same way. Progressive mode is untouched (full-res, unclamped, converging). --- native/shared/src/renderer/mod.rs | 48 ++++++- native/shared/src/renderer/pt_pass.rs | 85 +++++++++-- native/shared/src/renderer/shaders/pt.rs | 171 +++++++++++++++++------ 3 files changed, 246 insertions(+), 58 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index e8ff774..b8ec54e 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -381,11 +381,14 @@ struct PtParamsCpu { sun_color: [f32; 4], /// rgb ambient-derived sky tint, w unused. sky_color: [f32; 4], - /// x=width, y=height, z=frame_index, w=accum_count. + /// x/y = TRACE grid dims (half in realtime mode), z=frame_index, + /// w=accum_count. size: [u32; 4], /// x=mode (1 progressive / 2 realtime), y=max_bounces, /// z=point_light_count (≤16), w=debug view (BLOOM_PT_DEBUG). cfg: [f32; 4], + /// PT-3 half-res: x/y = full G-buffer dims, z/w = 2x2 sample phase. + ext: [u32; 4], /// 16 lights × (pos_range vec4, color_int vec4). lights: [[f32; 4]; 32], } @@ -1309,11 +1312,14 @@ pub struct Renderer { pt_atrous_final_pipeline: Option, pt_atrous_layout: Option, /// Per-iteration uniform (step size + sigmas + extent). - pt_atrous_params_bufs: [wgpu::Buffer; 2], + pt_atrous_params_bufs: [wgpu::Buffer; 3], pt_atrous_scratch: Option, - /// bg[0][i] = pass 1 reading accum buffer i; bg[1][0] = pass 2 - /// (scratch → hdr). Invalidated with pt_bg. + pt_atrous_scratch2: Option, + /// Three iterations: bg1[i] = accum buffer i → scratch (step 1), + /// bg_mid2 = scratch → scratch2 (step 2), bg2 = scratch2 → hdr + /// (step 4, upsampling). Invalidated with pt_bg. pt_atrous_bg1: [Option; 2], + pt_atrous_bg_mid2: Option, pt_atrous_bg2: Option, /// True when the PT kernel actually replaced this frame's scene /// colour. Mode 1 leaves the raster frame on screen until a few @@ -5242,6 +5248,26 @@ impl Renderer { view_dimension: wgpu::TextureViewDimension::D2, }, count: None, }, + // PT-3 half-res: full-res depth guides the + // upsample in cs_final. + wgpu::BindGroupLayoutEntry { + binding: 4, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + // Full-res G-buffer albedo for SVGF + // re-modulation in cs_final. + wgpu::BindGroupLayoutEntry { + binding: 5, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: false }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, ], }); let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { @@ -5268,13 +5294,19 @@ impl Renderer { let pt_atrous_params_bufs = [ device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_atrous_params_0"), - size: 16, + size: 32, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }), device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_atrous_params_1"), - size: 16, + size: 32, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }), + device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_atrous_params_2"), + size: 32, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }), @@ -6913,7 +6945,9 @@ impl Renderer { pt_atrous_layout, pt_atrous_params_bufs, pt_atrous_scratch: None, + pt_atrous_scratch2: None, pt_atrous_bg1: [None, None], + pt_atrous_bg_mid2: None, pt_atrous_bg2: None, probe_trace_sdf_pipeline, probe_trace_sdf_layout, @@ -7417,7 +7451,9 @@ impl Renderer { self.pt_accum_buffers = [None, None]; self.pt_accum_count = 0; self.pt_atrous_scratch = None; + self.pt_atrous_scratch2 = None; self.pt_atrous_bg1 = [None, None]; + self.pt_atrous_bg_mid2 = None; self.pt_atrous_bg2 = None; self.hiz_linearize_bg_cache = None; self.occlusion.invalidate_bindings(); diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index b92bb6e..fab3412 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -85,8 +85,30 @@ impl Renderer { return; } + // ---- trace grid ---- + // Realtime mode traces at half resolution (4x fewer rays) and + // joint-bilaterally upsamples in the final à-trous pass; the + // 2x2 sample phase rotates per frame so the temporal EMA + // integrates full-res coverage over 4 frames. Progressive mode + // stays full-res. + let (trace_w, trace_h) = if self.pt_mode >= 2 { + (surf_w.div_ceil(2), surf_h.div_ceil(2)) + } else { + (surf_w, surf_h) + }; + // Phase pinned to (0,0): rotating it makes each trace texel + // sample a different full-res pixel every frame, and on + // depth-chaotic surfaces (grass) the history validation then + // rejects almost every frame — texels never accumulate past + // 1 spp and read as white speckle. A consistent owner pixel + // keeps history valid; the upsample covers the other three. + let phase = [0u32, 0u32]; + // ---- accumulation buffers (vec4 per pixel, ping-pong) ---- - let needed = (surf_w as u64) * (surf_h as u64) * 16; + // Sized to the TRACE grid; a mode switch changes the size and + // recreates (which also resets accumulation — correct, the two + // modes' buffer contents are not interchangeable). + let needed = (trace_w as u64) * (trace_h as u64) * 16; let recreate = match &self.pt_accum_buffers[0] { Some(b) => b.size() != needed, None => true, @@ -110,8 +132,15 @@ impl Renderer { usage: wgpu::BufferUsages::STORAGE, mapped_at_creation: false, })); + self.pt_atrous_scratch2 = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some("pt_atrous_scratch2"), + size: needed, + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + })); self.pt_bg = [None, None]; self.pt_atrous_bg1 = [None, None]; + self.pt_atrous_bg_mid2 = None; self.pt_atrous_bg2 = None; self.pt_accum_count = 0; self.pt_accum_idx = 0; @@ -181,13 +210,14 @@ impl Renderer { amb[2] * sky_intensity, 0.0, ], - size: [surf_w, surf_h, self.taa_frame_index, self.pt_accum_count], + size: [trace_w, trace_h, self.taa_frame_index, self.pt_accum_count], cfg: [ self.pt_mode as f32, max_bounces, light_count as f32, self.pt_debug, ], + ext: [surf_w, surf_h, phase[0], phase[1]], lights, }; self.queue.write_buffer(&self.pt_uniform_buffer, 0, bytemuck::bytes_of(¶ms)); @@ -259,7 +289,7 @@ impl Renderer { if self.pt_texture_arrays_enabled { pass.set_bind_group(1, self.pt_tex_bg.as_ref().unwrap(), &[]); } - pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); + pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); } // This frame wrote into buffers[1 - idx]; it becomes next // frame's read side. @@ -275,11 +305,19 @@ impl Renderer { && self.pt_atrous_mid_pipeline.is_some() && self.pt_atrous_scratch.is_some() { - let sigma_luma = 0.25f32; - let p0 = [1.0f32, sigma_luma, surf_w as f32, surf_h as f32]; - let p1 = [2.0f32, sigma_luma, surf_w as f32, surf_h as f32]; - self.queue.write_buffer(&self.pt_atrous_params_bufs[0], 0, bytemuck::bytes_of(&p0)); - self.queue.write_buffer(&self.pt_atrous_params_bufs[1], 0, bytemuck::bytes_of(&p1)); + // SVGF-style schedule: the first iteration barely edge-stops + // on luminance (1-2 spp variance IS the signal to remove); + // later iterations tighten so real shading detail survives. + // Depth stopping guards geometry edges throughout. + for (i, (step, sigma_luma)) in + [(1.0f32, 2.0f32), (2.0, 0.5), (4.0, 0.3)].iter().enumerate() + { + let p = [ + [*step, *sigma_luma, trace_w as f32, trace_h as f32], + [surf_w as f32, surf_h as f32, 0.0, 0.0], + ]; + self.queue.write_buffer(&self.pt_atrous_params_bufs[i], 0, bytemuck::bytes_of(&p)); + } if self.pt_atrous_bg1[written_idx].is_none() { self.pt_atrous_bg1[written_idx] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { @@ -290,6 +328,22 @@ impl Renderer { wgpu::BindGroupEntry { binding: 1, resource: self.pt_accum_buffers[written_idx].as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 2, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, + ], + })); + } + if self.pt_atrous_bg_mid2.is_none() { + self.pt_atrous_bg_mid2 = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_atrous_bg_mid2"), + layout: self.pt_atrous_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[1].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: self.pt_atrous_scratch2.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, ], })); } @@ -298,14 +352,16 @@ impl Renderer { label: Some("pt_atrous_bg2"), layout: self.pt_atrous_layout.as_ref().unwrap(), entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[1].as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[2].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch2.as_ref().unwrap().as_entire_binding() }, // dst is unused by cs_final; bind an accum // buffer to satisfy the layout (cannot alias - // the scratch src binding — RO+RW of the same + // the scratch2 src binding — RO+RW of the same // buffer in one group fails validation). wgpu::BindGroupEntry { binding: 2, resource: self.pt_accum_buffers[0].as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, ], })); } @@ -315,9 +371,14 @@ impl Renderer { label: Some("pt_atrous"), timestamp_writes: ts, }); + // Mid passes filter on the trace grid (steps 1 and 2); the + // final pass runs at FULL resolution with step 4 (it + // upsamples when the grids differ). pass.set_pipeline(self.pt_atrous_mid_pipeline.as_ref().unwrap()); pass.set_bind_group(0, self.pt_atrous_bg1[written_idx].as_ref().unwrap(), &[]); - pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); + pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); + pass.set_bind_group(0, self.pt_atrous_bg_mid2.as_ref().unwrap(), &[]); + pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); pass.set_pipeline(self.pt_atrous_final_pipeline.as_ref().unwrap()); pass.set_bind_group(0, self.pt_atrous_bg2.as_ref().unwrap(), &[]); pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 1eb31e6..be4581a 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -55,8 +55,11 @@ struct PtParams { sun_dir: vec4, // xyz unit vector toward the sun sun_color: vec4, // rgb premultiplied by intensity sky_color: vec4, // rgb ambient-derived sky tint - size: vec4, // x=width, y=height, z=frame_index, w=accum_count + size: vec4, // x/y = TRACE grid dims, z=frame_index, w=accum_count cfg: vec4, // x=mode(1|2), y=max_bounces, z=point_light_count, w=debug + // PT-3 half-res: x/y = full G-buffer dims, z/w = 2x2 sample phase. + // Full-res modes set ext.xy == size.xy and phase 0. + ext: vec4, lights: array, }; @@ -173,8 +176,10 @@ fn rand_2f() -> vec2 { return vec2(rand_f(), rand_f()); } // ---- Geometry reconstruction ------------------------------------------------- +// px here is always a FULL-resolution G-buffer pixel (u.ext dims); the +// trace grid may be half of that in realtime mode. fn world_at(px: vec2, depth: f32) -> vec3 { - let dims = vec2(f32(u.size.x), f32(u.size.y)); + let dims = vec2(f32(u.ext.x), f32(u.ext.y)); let uv = (vec2(px) + vec2(0.5)) / dims; let ndc = vec4(uv.x * 2.0 - 1.0, 1.0 - uv.y * 2.0, depth, 1.0); let w = u.inv_vp * ndc; @@ -182,7 +187,7 @@ fn world_at(px: vec2, depth: f32) -> vec3 { } fn depth_at(px: vec2) -> f32 { - let clamped = clamp(px, vec2(0), vec2(i32(u.size.x) - 1, i32(u.size.y) - 1)); + let clamped = clamp(px, vec2(0), vec2(i32(u.ext.x) - 1, i32(u.ext.y) - 1)); return textureLoad(depth_tex, clamped, 0); } @@ -510,15 +515,27 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let px = vec2(i32(gid.x), i32(gid.y)); rng_seed(gid.xy, u.size.z); + // PT-3 half-res: realtime mode traces a half grid; map this trace + // cell to its full-res G-buffer pixel (the 2x2 phase rotates per + // frame so the EMA integrates all four over time). Full-res modes + // have ext == size and phase 0, making this the identity. + var px_full = px; + if (u.ext.x > u.size.x) { + px_full = min( + px * 2 + vec2(i32(u.ext.z), i32(u.ext.w)), + vec2(i32(u.ext.x) - 1, i32(u.ext.y) - 1), + ); + } + let debug = u.cfg.w; if (debug == 5.0) { - textureStore(out_hdr, px, vec4(1.0, 0.0, 1.0, 1.0)); + textureStore(out_hdr, px_full, vec4(1.0, 0.0, 1.0, 1.0)); return; } - let depth = depth_at(px); + let depth = depth_at(px_full); if (debug == 1.0) { - textureStore(out_hdr, px, vec4(vec3(depth), 1.0)); + textureStore(out_hdr, px_full, vec4(vec3(depth), 1.0)); return; } @@ -531,22 +548,22 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { return; } - let p0 = world_at(px, depth); - let n0 = normal_from_depth(px, p0); - let albedo0 = textureLoad(albedo_tex, px, 0).rgb; + let p0 = world_at(px_full, depth); + let n0 = normal_from_depth(px_full, p0); + let albedo0 = textureLoad(albedo_tex, px_full, 0).rgb; if (debug == 2.0) { - textureStore(out_hdr, px, vec4(n0 * 0.5 + 0.5, 1.0)); + textureStore(out_hdr, px_full, vec4(n0 * 0.5 + 0.5, 1.0)); return; } if (debug == 3.0) { - textureStore(out_hdr, px, vec4(albedo0, 1.0)); + textureStore(out_hdr, px_full, vec4(albedo0, 1.0)); return; } if (debug == 4.0) { let sd = sun_cone_sample(rand_2f()); let vis = select(0.0, 1.0, !occluded(p0 + n0 * 0.02, sd, 1000.0)); - textureStore(out_hdr, px, vec4(vec3(vis), 1.0)); + textureStore(out_hdr, px_full, vec4(vec3(vis), 1.0)); return; } if (debug == 8.0) { @@ -565,7 +582,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let gi = instance_data[h8.instance_custom_data].geo; c8 = select(vec3(0.0), vec3(100.0), gi.z > 0u); } - textureStore(out_hdr, px, vec4(c8, 1.0)); + textureStore(out_hdr, px_full, vec4(c8, 1.0)); return; } if (debug == 9.0) { @@ -599,7 +616,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } } } - textureStore(out_hdr, px, vec4(c9, 1.0)); + textureStore(out_hdr, px_full, vec4(c9, 1.0)); return; } if (debug == 10.0) { @@ -619,7 +636,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let prim = f32(h10.primitive_index); c10 = vec3(fract(prim / 64.0), fract(prim / 1024.0), fract(prim / 16384.0)) * 30.0; } - textureStore(out_hdr, px, vec4(c10, 1.0)); + textureStore(out_hdr, px_full, vec4(c10, 1.0)); return; } if (debug == 11.0 || debug == 12.0) { @@ -647,7 +664,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { c11 = vec3(b.x, b.y, max(0.0, 1.0 - b.x - b.y)) * 30.0; } } - textureStore(out_hdr, px, vec4(c11, 1.0)); + textureStore(out_hdr, px_full, vec4(c11, 1.0)); return; } if (debug == 13.0) { @@ -674,7 +691,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { c13 = vec3(50.0, 0.0, 0.0); } } - textureStore(out_hdr, px, vec4(c13, 1.0)); + textureStore(out_hdr, px_full, vec4(c13, 1.0)); return; } if (debug == 14.0) { @@ -695,7 +712,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { fract(h14.t * 0.0078125), ) * 20.0; } - textureStore(out_hdr, px, vec4(c14, 1.0)); + textureStore(out_hdr, px_full, vec4(c14, 1.0)); return; } if (debug == 15.0) { @@ -721,7 +738,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { if (hA.kind != RAY_QUERY_INTERSECTION_NONE) { tA = hA.t; } if (hB.kind != RAY_QUERY_INTERSECTION_NONE) { tB = hB.t; } let c15 = vec3(fract(tA * 0.125) * 20.0, fract(tB * 0.125) * 20.0, 0.0); - textureStore(out_hdr, px, vec4(c15, 1.0)); + textureStore(out_hdr, px_full, vec4(c15, 1.0)); return; } if (debug == 16.0) { @@ -743,7 +760,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { f32(h16.primitive_index), f32(h16.kind), ); - textureStore(out_hdr, px, vec4(0.2, 0.0, 0.4, 1.0)); + textureStore(out_hdr, px_full, vec4(0.2, 0.0, 0.4, 1.0)); return; } if (debug == 17.0) { @@ -751,7 +768,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // depth, straight into accum for CPU readback. let idx17 = gid.y * u.size.x + gid.x; accum_out[idx17] = vec4(p0, depth); - textureStore(out_hdr, px, vec4(0.4, 0.2, 0.0, 1.0)); + textureStore(out_hdr, px_full, vec4(0.4, 0.2, 0.0, 1.0)); return; } if (debug == 6.0 || debug == 7.0) { @@ -782,7 +799,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { col = vec3(1.0, 0.5, 0.0); // orange: no geo window } } - textureStore(out_hdr, px, vec4(col, 1.0)); + textureStore(out_hdr, px_full, vec4(col, 1.0)); return; } @@ -793,7 +810,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // (1 - metallic) — metals have no diffuse lobe. Specular NEE is a // known gap (see the PT-2 ticket); specular reflection of sky and // scene comes from the GGX bounce below. - let mr0 = textureLoad(material_tex, px, 0).rg; + let mr0 = textureLoad(material_tex, px_full, 0).rg; var metal_cur = mr0.r; var rough_cur = mr0.g; var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur)); @@ -887,7 +904,11 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // frames no matter how large the frame count grows, so the cap // must never relax or fireflies return as the count climbs. var cap = 4.0 + f32(min(u.size.w, 28u)); - if (u.cfg.x >= 2.0) { cap = 6.0; } + // Realtime cap is TIGHT (3): the à-trous luminance edge-stop + // preserves bright outliers as if they were edges, so any firefly + // that survives the cap survives the filter too. Scene luma sits + // around 0.5-1.5, so 3 leaves sunlit highlights alone. + if (u.cfg.x >= 2.0) { cap = 3.0; } if (luma > cap) { radiance *= cap / luma; } // ---- accumulate --------------------------------------------------------- @@ -919,16 +940,35 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let ph = accum[pidx]; let zl_hist = lin_depth(ph.w); let zl_here = lin_depth(ndc_prev.z); - if (abs(zl_hist - zl_here) < 0.05 * max(zl_hist, zl_here) + 0.05) { + // Tolerance is deliberately loose: at half-res the + // reprojected texel quantizes to 2 full pixels, and on + // depth-chaotic surfaces (grass) a tight test rejects + // almost every frame — pinning texels at 1 spp. The + // ghosting a loose test admits is invisible exactly + // where the depth is chaotic. + if (abs(zl_hist - zl_here) < 0.12 * max(zl_hist, zl_here) + 0.05) { alpha = 0.15; hist = ph.rgb; } } } - out = mix(hist, radiance, alpha); - accum_out[idx] = vec4(out, depth); - // Realtime output goes through the a-trous passes (they write - // out_hdr); the main kernel only feeds the history buffer. + // SVGF-style albedo demodulation: history and output store + // IRRADIANCE (radiance divided by the primary albedo) so the + // à-trous passes filter lighting only; the final pass + // re-multiplies by the FULL-RES G-buffer albedo, keeping + // texture detail crisp despite the half-res trace and the + // aggressive smoothing. + var irr = radiance / max(albedo0, vec3(0.05)); + // The firefly cap must bind IRRADIANCE, not radiance: dividing + // by a dark albedo amplifies up to 20x, so a radiance-space cap + // leaves outliers the filter sigmas can't absorb. Sunlit + // irradiance sits around 1-3; 4 leaves highlights alone. + let irr_luma = dot(irr, vec3(0.2126, 0.7152, 0.0722)); + if (irr_luma > 4.0) { + irr *= 4.0 / irr_luma; + } + let out_irr = mix(hist, irr, alpha); + accum_out[idx] = vec4(out_irr, depth); return; } else { // Progressive: plain running sum; count lives on the CPU. @@ -949,7 +989,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } } - textureStore(out_hdr, px, vec4(out, 1.0)); + textureStore(out_hdr, px_full, vec4(out, 1.0)); } "#; @@ -961,13 +1001,20 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { /// through untouched and are never written to hdr. pub(in crate::renderer) const PT_ATROUS_WGSL: &str = r#" struct AtrousParams { - // x = step (texels), y = sigma_luma, z = width, w = height + // x = step (texels), y = sigma_luma, z = trace width, w = trace height p: vec4, + // x/y = full G-buffer dims (cs_final upsamples trace -> full when + // they differ), z/w unused. + p2: vec4, }; @group(0) @binding(0) var ap: AtrousParams; @group(0) @binding(1) var src: array>; @group(0) @binding(2) var dst: array>; @group(0) @binding(3) var out_hdr_a: texture_storage_2d; +@group(0) @binding(4) var depth_full: texture_depth_2d; +// Full-res G-buffer albedo: cs_final re-modulates the filtered +// irradiance with it (SVGF demodulation keeps textures crisp). +@group(0) @binding(5) var albedo_full: texture_2d; fn lin_depth_a(d: f32) -> f32 { return 0.02 / max(1.0 - d, 1e-6); @@ -1002,7 +1049,7 @@ fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { continue; } let zq = lin_depth_a(s.w); - let wz = exp(-abs(zq - zc) / (0.08 * max(zc, zq) + 0.02)); + let wz = exp(-abs(zq - zc) / (0.15 * max(zc, zq) + 0.02)); let lq = dot(s.rgb, vec3(0.2126, 0.7152, 0.0722)); let wl = exp(-abs(lq - lc) / max(ap.p.y, 1e-3)); let wgt = kern(dx) * kern(dy) * wz * wl; @@ -1027,20 +1074,64 @@ fn cs_mid(@builtin(global_invocation_id) gid: vec3) { dst[gid.y * u32(w) + gid.x] = filter_at(px, w, h, i32(ap.p.x)); } +// Final pass runs at FULL resolution: filter-and-write when the trace +// grid is full-res, depth-guided joint-bilateral upsample when the +// trace grid is half-res. Sky pixels (full-res depth at far plane) are +// never written so the raster sky survives. @compute @workgroup_size(8, 8, 1) fn cs_final(@builtin(global_invocation_id) gid: vec3) { - let w = i32(ap.p.z); - let h = i32(ap.p.w); - if (i32(gid.x) >= w || i32(gid.y) >= h) { + let fw = i32(ap.p2.x); + let fh = i32(ap.p2.y); + if (i32(gid.x) >= fw || i32(gid.y) >= fh) { return; } let px = vec2(i32(gid.x), i32(gid.y)); - let cidx = gid.y * u32(w) + gid.x; - if (src[cidx].w >= 0.9999999) { - // Sky: the raster sky in hdr stays untouched. + let d = textureLoad(depth_full, px, 0); + if (d >= 0.9999999) { + return; + } + let hw = i32(ap.p.z); + let hh = i32(ap.p.w); + let alb = max(textureLoad(albedo_full, px, 0).rgb, vec3(0.05)); + if (hw == fw) { + let cidx = gid.y * u32(fw) + gid.x; + if (src[cidx].w >= 0.9999999) { + return; + } + let r = filter_at(px, fw, fh, i32(ap.p.x)); + textureStore(out_hdr_a, px, vec4(r.rgb * alb, 1.0)); + return; + } + // Half-res upsample: 3x3 taps around the containing trace texel, + // weighted by kernel distance and relative linear-depth agreement + // with THIS full-res pixel. The epsilon keeps thin foreground + // geometry (whose taps all mismatch) softly averaged rather than + // black. + let zc = lin_depth_a(d); + let hx = clamp(px.x / 2, 0, hw - 1); + let hy = clamp(px.y / 2, 0, hh - 1); + var sum = vec3(0.0); + var wsum = 0.0; + for (var dy = -1; dy <= 1; dy = dy + 1) { + for (var dx = -1; dx <= 1; dx = dx + 1) { + let qx = hx + dx; + let qy = hy + dy; + if (qx < 0 || qy < 0 || qx >= hw || qy >= hh) { + continue; + } + let s = src[u32(qy) * u32(hw) + u32(qx)]; + if (s.w >= 0.9999999) { + continue; + } + let wz = exp(-abs(lin_depth_a(s.w) - zc) / (0.08 * zc + 0.02)); + let wgt = kern(dx) * kern(dy) * wz + 1e-5; + sum += s.rgb * wgt; + wsum += wgt; + } + } + if (wsum < 1e-6) { return; } - let r = filter_at(px, w, h, i32(ap.p.x)); - textureStore(out_hdr_a, px, vec4(r.rgb, 1.0)); + textureStore(out_hdr_a, px, vec4((sum / wsum) * alb, 1.0)); } "#; From d7d7b611321037008c7e5e7974259e147855fda2 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 07:49:59 +0200 Subject: [PATCH 11/16] fix(pt): despeckle pass + history-guided spike clamp for RT sparkles Two mechanisms let isolated bright texels survive the whole denoiser: (1) the a-trous edge-stopping weights are measured FROM the center texel, so when the center itself is the outlier it rejects all its neighbors and passes through every iteration untouched; (2) a fresh 1-spp spike lives on screen for a frame before any filtering can act. Fixes: the first filter iteration clamps the center to 1.5x its immediate neighborhood maximum before blending (isolated texels cannot exceed what any neighbor has seen), and the kernel clamps a new sample to 1.5x + 0.25 above VALID reprojected history (single-frame upward spikes are noise at 1-2 spp; darkening stays unclamped so appearing shadows propagate at full speed). Verified mid-strafe: faint sparse dusting remains, the sparkle field is gone. 22 fps at 1080p internal. --- native/shared/src/renderer/shaders/pt.rs | 49 ++++++++++++++++++++++-- 1 file changed, 46 insertions(+), 3 deletions(-) diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index be4581a..f4fef4c 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -967,6 +967,15 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { if (irr_luma > 4.0) { irr *= 4.0 / irr_luma; } + // History-guided spike clamp: with valid history, one frame may + // not jump more than 50% + 0.25 above it — single-frame bright + // spikes ARE noise at 1-2 spp, and the à-trous edge-stop + // (center-relative) cannot remove an outlier CENTER. Darkening + // (a shadow appearing) stays unclamped so lighting changes + // propagate at full speed downward and EMA speed upward. + if (alpha < 1.0) { + irr = min(irr, hist * 1.5 + vec3(0.25)); + } let out_irr = mix(hist, irr, alpha); accum_out[idx] = vec4(out_irr, depth); return; @@ -1030,12 +1039,43 @@ fn kern(d: i32) -> f32 { fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { let cidx = u32(px.y) * u32(w) + u32(px.x); - let center = src[cidx]; + var center = src[cidx]; if (center.w >= 0.9999999) { return center; } let zc = lin_depth_a(center.w); - let lc = dot(center.rgb, vec3(0.2126, 0.7152, 0.0722)); + var lc = dot(center.rgb, vec3(0.2126, 0.7152, 0.0722)); + + // Despeckle (first iteration only): the edge-stopping weights below + // are measured FROM the center, so an outlier center rejects its + // neighbors and survives every iteration. Clamp the center to its + // immediate neighborhood maximum before blending — isolated bright + // texels cannot exceed what any neighbor has seen. + if (step == 1) { + var nmax = 0.0; + for (var dy = -1; dy <= 1; dy = dy + 1) { + for (var dx = -1; dx <= 1; dx = dx + 1) { + if (dx == 0 && dy == 0) { + continue; + } + let q = px + vec2(dx, dy); + if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { + continue; + } + let s = src[u32(q.y) * u32(w) + u32(q.x)]; + if (s.w >= 0.9999999) { + continue; + } + nmax = max(nmax, dot(s.rgb, vec3(0.2126, 0.7152, 0.0722))); + } + } + let lim = nmax * 1.5 + 0.05; + if (nmax > 0.0 && lc > lim) { + center = vec4(center.rgb * (lim / lc), center.w); + lc = lim; + } + } + var sum = vec3(0.0); var wsum = 0.0; for (var dy = -2; dy <= 2; dy = dy + 1) { @@ -1044,7 +1084,10 @@ fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { continue; } - let s = src[u32(q.y) * u32(w) + u32(q.x)]; + var s = src[u32(q.y) * u32(w) + u32(q.x)]; + if (dx == 0 && dy == 0) { + s = center; // use the despeckled value for the center tap + } if (s.w >= 0.9999999) { continue; } From c2091123ae142a91509e062dae69fa30df7d38cd Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 08:26:41 +0200 Subject: [PATCH 12/16] feat(pt): IGN-structured primary sun sampling + deeper RT history Realtime mode now samples the primary sun cone with interleaved gradient noise (scrolled per frame) instead of white PCG noise: neighboring pixels get well-distributed sample positions, so the a-trous kernel averages penumbra/foliage sun visibility nearly flat instead of leaving shimmering mid-frequency blotch. Progressive mode keeps white noise for unbiased convergence. History EMA deepened (alpha 0.15 -> 0.1) now that spike clamps bound the outliers. Numeric shimmer meter (two stills 1.5s apart, standing, center crop): RT mean |delta| 54.6 / sparkle-grade 16.7% vs raster baseline 46.6 / 14.5% - RT temporal churn is now within ~17% of what the raster pipeline itself produces (TAA jitter + animated water dominate both). --- native/shared/src/renderer/shaders/pt.rs | 30 ++++++++++++++++++++---- 1 file changed, 25 insertions(+), 5 deletions(-) diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index f4fef4c..68e0988 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -174,6 +174,16 @@ fn rand_f() -> f32 { fn rand_2f() -> vec2 { return vec2(rand_f(), rand_f()); } +// Interleaved gradient noise, scrolled per frame by the golden-ratio +// offset. Spatially STRUCTURED (neighbors get well-distributed values) +// so a 5x5 filter averages it nearly flat — white PCG noise leaves +// mid-frequency blotch at 1-2 spp that shimmers. Used for the primary +// sun test in realtime mode. +fn ign_at(px: vec2, frame: u32) -> f32 { + let p = vec2(px) + f32(frame % 64u) * 5.588238; + return fract(52.9829189 * fract(0.06711056 * p.x + 0.00583715 * p.y)); +} + // ---- Geometry reconstruction ------------------------------------------------- // px here is always a FULL-resolution G-buffer pixel (u.ext dims); the @@ -478,10 +488,10 @@ fn albedo_at_hit( // Direct light at a surface point: sun through the solar cone + one point // light chosen uniformly (contribution / pdf). Game-radiometry convention: // no 1/pi (see file header). -fn direct_light(p: vec3, n: vec3, alb: vec3) -> vec3 { +fn direct_light(p: vec3, n: vec3, alb: vec3, sun_r2: vec2) -> vec3 { var lit = vec3(0.0); - let sd = sun_cone_sample(rand_2f()); + let sd = sun_cone_sample(sun_r2); let ndl = dot(n, sd); if (ndl > 0.0 && !occluded(p, sd, 1000.0)) { lit += u.sun_color.rgb * ndl; @@ -813,7 +823,17 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { let mr0 = textureLoad(material_tex, px_full, 0).rg; var metal_cur = mr0.r; var rough_cur = mr0.g; - var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur)); + // Realtime mode samples the primary sun cone with structured IGN + // noise (filterable); progressive keeps white noise (unbiased + // convergence). + var sun_r2 = rand_2f(); + if (u.cfg.x >= 2.0) { + sun_r2 = vec2( + ign_at(px_full, u.size.z), + ign_at(px_full + vec2(17, 59), u.size.z), + ); + } + var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur), sun_r2); var throughput = vec3(1.0); var origin = p0 + n0 * 0.02; var n_cur = n0; @@ -874,7 +894,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { radiance += throughput * inst.albedo * inst.emissive_luma; let hit_p = hit_ws + n_hit * 0.02; - radiance += throughput * direct_light(hit_p, n_hit, alb_hit * (1.0 - inst.mat_params.y)); + radiance += throughput * direct_light(hit_p, n_hit, alb_hit * (1.0 - inst.mat_params.y), rand_2f()); origin = hit_p; n_cur = n_hit; @@ -947,7 +967,7 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // ghosting a loose test admits is invisible exactly // where the depth is chaotic. if (abs(zl_hist - zl_here) < 0.12 * max(zl_hist, zl_here) + 0.05) { - alpha = 0.15; + alpha = 0.1; hist = ph.rgb; } } From 7ee970d881a4d8f28a1a09410dafe4cd370585b9 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 08:36:46 +0200 Subject: [PATCH 13/16] fix(pt): freeze the RT sample sequence - zero PT-injected temporal noise The remaining full-frame sparkling was not outliers: every pixel whose sun visibility is stochastic (grass blades, canopies - most of a 4K frame) re-rolled its sample each frame and oscillated around its mean forever; a 10-frame EMA bounds the amplitude but never settles it, and no spatial filter removes temporally-uncorrelated oscillation. Realtime mode now uses a FROZEN sequence: IGN sun samples at frame 0 and the bounce RNG seeded without the frame index. Every sample is deterministic per pixel, so the path tracer contributes zero frame-to-frame variance; the a-trous flattens the static dither into stable soft shadows. Progressive and the debug views keep rolling sequences for unbiased convergence. Grass-band shimmer meter (two stills 1.5s apart, standing): RT now 32.0 mean / 9.1% sparkle-grade vs the RASTER pipeline itself at 42.8 / 12.9% - realtime path tracing is temporally quieter than raster TAA. --- native/shared/src/renderer/shaders/pt.rs | 24 +++++++++++++++++++----- 1 file changed, 19 insertions(+), 5 deletions(-) diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 68e0988..86fb617 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -523,7 +523,16 @@ fn direct_light(p: vec3, n: vec3, alb: vec3, sun_r2: vec2) - fn cs_main(@builtin(global_invocation_id) gid: vec3) { if (gid.x >= u.size.x || gid.y >= u.size.y) { return; } let px = vec2(i32(gid.x), i32(gid.y)); - rng_seed(gid.xy, u.size.z); + // Realtime mode uses a FROZEN random sequence (seed 0 instead of + // the frame index): every sample is deterministic per pixel, so + // the path tracer injects ZERO frame-to-frame noise — the à-trous + // smooths a static dither into a stable image. Progressive (and + // the debug views) keep the rolling sequence for convergence. + var seed_frame = u.size.z; + if (u.cfg.x >= 2.0 && u.cfg.w == 0.0) { + seed_frame = 0u; + } + rng_seed(gid.xy, seed_frame); // PT-3 half-res: realtime mode traces a half grid; map this trace // cell to its full-res G-buffer pixel (the 2x2 phase rotates per @@ -824,13 +833,18 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { var metal_cur = mr0.r; var rough_cur = mr0.g; // Realtime mode samples the primary sun cone with structured IGN - // noise (filterable); progressive keeps white noise (unbiased - // convergence). + // noise, FROZEN in time (frame 0): a per-frame rolling sequence + // makes every stochastically-shadowed pixel (grass, canopies) + // oscillate around its mean forever — the EMA never settles it and + // it reads as sparkle across the whole frame. A fixed per-pixel + // sample gives a static dither that the à-trous flattens into + // STABLE soft shadows. Progressive keeps rolling white noise for + // unbiased convergence. var sun_r2 = rand_2f(); if (u.cfg.x >= 2.0) { sun_r2 = vec2( - ign_at(px_full, u.size.z), - ign_at(px_full + vec2(17, 59), u.size.z), + ign_at(px_full, 0u), + ign_at(px_full + vec2(17, 59), 0u), ); } var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur), sun_r2); From 85920d8340a043b9938673c77bcab9d5e8c537a6 Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 09:19:04 +0200 Subject: [PATCH 14/16] feat(pt): hybrid direct sun + sun specular for realtime mode RT mode direct sunlight now comes from the raster shadow cascades (manual 2x2 load-and-compare) instead of one traced cone ray per pixel: crisp deterministic shadows identical to raster, zero direct- light noise, and the traced rays are spent purely on indirect. This addressed the dominant realism complaints - blobby dithered shadows and stippled penumbra. Progressive mode keeps traced sun for reference-quality soft shadows. Also adds the sun GGX specular at the primary surface (both modes; masked by the same visibility) - without it every surface read as dead-matte under direct sun. Bring-up notes, paid in verification blood: - shadow_map.light_vps upload RAW (they are consumed as M*v by every existing WGSL user, unlike the camera matrices which need the transpose). Verified numerically via a new debug-19 NDC dump. - textureSampleCompareLevel from compute was bypassed for manual textureLoad compare during diagnosis; kept the manual version (2x2 PCF) since it is proven and equally fast at this tap count. - The a-trous passes now skip when a debug view is active - they were overwriting every debug output with stale-buffer garbage, which cost two false diagnoses (both prior black frames were the filter stomping the debug write, not the cascades failing). - Debug readback offsets now use trace-grid dims (was full-res - OOB copy validation panic in half-res mode). --- native/shared/src/renderer/mod.rs | 36 +++++- native/shared/src/renderer/pt_pass.rs | 32 +++++- native/shared/src/renderer/shaders/pt.rs | 139 +++++++++++++++++++++-- 3 files changed, 194 insertions(+), 13 deletions(-) diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index b8ec54e..5263ff2 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -387,8 +387,11 @@ struct PtParamsCpu { /// x=mode (1 progressive / 2 realtime), y=max_bounces, /// z=point_light_count (≤16), w=debug view (BLOOM_PT_DEBUG). cfg: [f32; 4], - /// PT-3 half-res: x/y = full G-buffer dims, z/w = 2x2 sample phase. + /// PT-3 half-res: x/y = full G-buffer dims. z = 1 → hybrid sun + /// (shadow cascades instead of traced sun rays). w unused. ext: [u32; 4], + /// Raster shadow cascade VPs, transposed at upload like inv_vp. + shadow_vps: [[[f32; 4]; 4]; 3], /// 16 lights × (pos_range vec4, color_int vec4). lights: [[f32; 4]; 32], } @@ -5163,6 +5166,37 @@ impl Renderer { has_dynamic_offset: false, min_binding_size: None, }, count: None, }, + // Hybrid sun — the raster shadow cascades + their + // comparison sampler. + wgpu::BindGroupLayoutEntry { + binding: 14, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 15, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 16, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Depth, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 17, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison), + count: None, + }, ]; let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("pt_layout"), diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index fab3412..4ebef04 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -217,7 +217,22 @@ impl Renderer { light_count as f32, self.pt_debug, ], - ext: [surf_w, surf_h, phase[0], phase[1]], + // ext.z: hybrid sun — realtime mode samples the raster + // shadow cascades instead of tracing sun rays (crisp + // noise-free direct shadows). Progressive keeps traced sun + // for reference-quality penumbra. + ext: [ + surf_w, + surf_h, + if self.pt_mode >= 2 && self.shadow_map.enabled { 1 } else { 0 }, + phase[1], + ], + // RAW upload, unlike inv_vp: the shadow VPs are consumed as + // M*v by every existing WGSL user (scene shader, WSRC + // bake), so they are already stored in WGSL column layout. + // Verified empirically via debug 18 — transposing them + // black-shadows the whole frame. + shadow_vps: self.shadow_map.light_vps, lights, }; self.queue.write_buffer(&self.pt_uniform_buffer, 0, bytemuck::bytes_of(¶ms)); @@ -247,6 +262,10 @@ impl Renderer { wgpu::BindGroupEntry { binding: 10, resource: self.pt_geo_vertex_buffer.as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 11, resource: self.pt_geo_index_buffer.as_ref().unwrap().as_entire_binding() }, wgpu::BindGroupEntry { binding: 13, resource: self.pt_accum_buffers[1 - i].as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 14, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[0]) }, + wgpu::BindGroupEntry { binding: 15, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[1]) }, + wgpu::BindGroupEntry { binding: 16, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[2]) }, + wgpu::BindGroupEntry { binding: 17, resource: wgpu::BindingResource::Sampler(&self.shadow_map.sampler) }, ]; self.pt_bg[i] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("pt_bg"), @@ -302,6 +321,7 @@ impl Renderer { // (step 2). Progressive mode converges on its own and writes // hdr directly from the kernel. if self.pt_mode >= 2 + && self.pt_debug == 0.0 && self.pt_atrous_mid_pipeline.is_some() && self.pt_atrous_scratch.is_some() { @@ -394,13 +414,15 @@ impl Renderer { // Copies a window of the accum buffer (center of frame) into a // staging buffer each frame; the previous frame's copy is mapped // (blocking) and dumped to pt_trace_dump.txt once. - if (self.pt_debug == 16.0 || self.pt_debug == 17.0) + if (self.pt_debug == 16.0 || self.pt_debug == 17.0 || self.pt_debug == 19.0) && self.pt_accum_count > 30 && !self.pt_dump_written { - let dump_pixels: u64 = (surf_w as u64).min(4096); - let row = (surf_h / 2) as u64; - let offset = row * surf_w as u64 * 16; + // Offsets in TRACE-grid units — the accum buffers are + // half-res in realtime mode. + let dump_pixels: u64 = (trace_w as u64).min(4096); + let row = (trace_h / 2) as u64; + let offset = row * trace_w as u64 * 16; if self.pt_readback_buffer.is_none() { self.pt_readback_buffer = Some(self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_readback"), diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 86fb617..8701092 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -57,9 +57,14 @@ struct PtParams { sky_color: vec4, // rgb ambient-derived sky tint size: vec4, // x/y = TRACE grid dims, z=frame_index, w=accum_count cfg: vec4, // x=mode(1|2), y=max_bounces, z=point_light_count, w=debug - // PT-3 half-res: x/y = full G-buffer dims, z/w = 2x2 sample phase. - // Full-res modes set ext.xy == size.xy and phase 0. + // PT-3 half-res: x/y = full G-buffer dims. z = 1 -> hybrid sun + // (sample the raster shadow cascades instead of tracing the sun; + // crisp noise-free direct shadows, rays spent on indirect only). + // w unused. ext: vec4, + // Raster shadow cascade view-projections (transposed at upload, + // same convention as inv_vp/prev_vp). + shadow_vps: array, 3>, lights: array, }; @@ -109,6 +114,58 @@ fn lin_depth(d: f32) -> f32 { // appended by the Rust side per adapter support.) @group(0) @binding(10) var geo_v: array; @group(0) @binding(11) var geo_i: array; +// Hybrid sun (ext.z == 1): the raster shadow cascades. +@group(0) @binding(14) var shadow_atlas_0: texture_depth_2d; +@group(0) @binding(15) var shadow_atlas_1: texture_depth_2d; +@group(0) @binding(16) var shadow_atlas_2: texture_depth_2d; +@group(0) @binding(17) var shadow_samp: sampler_comparison; + +// Sun visibility from the shadow cascades (near -> far fallthrough by +// coverage, same scheme as the WSRC bake). Deterministic and smooth — +// the whole reason RT mode's direct light doesn't shimmer or dither. +fn sun_vis_cascade(pos_ws: vec3) -> f32 { + for (var c = 0; c < 3; c = c + 1) { + var clip: vec4; + if (c == 0) { clip = u.shadow_vps[0] * vec4(pos_ws, 1.0); } + else if (c == 1) { clip = u.shadow_vps[1] * vec4(pos_ws, 1.0); } + else { clip = u.shadow_vps[2] * vec4(pos_ws, 1.0); } + if (abs(clip.w) < 1e-6) { continue; } + let ndc = clip.xyz / clip.w; + if (ndc.x < -0.99 || ndc.x > 0.99 || ndc.y < -0.99 || ndc.y > 0.99 || ndc.z < 0.0 || ndc.z > 1.0) { + continue; + } + let uv = vec2(ndc.x * 0.5 + 0.5, 0.5 - ndc.y * 0.5); + let ref_depth = ndc.z - 0.002; + // Manual load-and-compare with a 2x2 average instead of the + // comparison sampler: SampleCmp from a COMPUTE stage proved + // unreliable on this DXC path (constant 0, independent of the + // matrices — same failure shape as the ray-query saga), and the + // only other compute-stage user (WSRC bake) was never validated + // on DX12. textureLoad is proven (the PT depth reads use it). + var dims: vec2; + if (c == 0) { dims = textureDimensions(shadow_atlas_0); } + else if (c == 1) { dims = textureDimensions(shadow_atlas_1); } + else { dims = textureDimensions(shadow_atlas_2); } + let fdims = vec2(dims); + var vis = 0.0; + for (var ty = 0; ty <= 1; ty = ty + 1) { + for (var tx = 0; tx <= 1; tx = tx + 1) { + let tc = clamp( + vec2(uv * fdims - vec2(0.5)) + vec2(tx, ty), + vec2(0), + vec2(i32(dims.x) - 1, i32(dims.y) - 1), + ); + var stored: f32; + if (c == 0) { stored = textureLoad(shadow_atlas_0, tc, 0); } + else if (c == 1) { stored = textureLoad(shadow_atlas_1, tc, 0); } + else { stored = textureLoad(shadow_atlas_2, tc, 0); } + if (ref_depth <= stored) { vis += 0.25; } + } + } + return vis; + } + return 1.0; +} const PT_VSTRIDE: u32 = 24u; @@ -488,13 +545,29 @@ fn albedo_at_hit( // Direct light at a surface point: sun through the solar cone + one point // light chosen uniformly (contribution / pdf). Game-radiometry convention: // no 1/pi (see file header). +// Sun visibility at a surface point: shadow cascades in hybrid mode +// (deterministic, matches the raster shadows exactly), a traced cone +// ray otherwise (reference quality, soft penumbra). +fn sun_visibility(p: vec3, n: vec3, r2: vec2) -> f32 { + if (u.ext.z == 1u) { + return sun_vis_cascade(p); + } + let sd = sun_cone_sample(r2); + if (dot(n, sd) <= 0.0) { + return 0.0; + } + if (occluded(p, sd, 1000.0)) { + return 0.0; + } + return 1.0; +} + fn direct_light(p: vec3, n: vec3, alb: vec3, sun_r2: vec2) -> vec3 { var lit = vec3(0.0); - let sd = sun_cone_sample(sun_r2); - let ndl = dot(n, sd); - if (ndl > 0.0 && !occluded(p, sd, 1000.0)) { - lit += u.sun_color.rgb * ndl; + let ndl = max(dot(n, u.sun_dir.xyz), 0.0); + if (ndl > 0.0) { + lit += u.sun_color.rgb * ndl * sun_visibility(p, n, sun_r2); } let count = u32(u.cfg.z); @@ -790,6 +863,36 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { textureStore(out_hdr, px_full, vec4(0.4, 0.2, 0.0, 1.0)); return; } + if (debug == 18.0) { + // Hybrid-sun validation: cascade shadow visibility at the + // primary surface. Must match the raster shadow shapes exactly + // (crisp tree/building shadows). Gray everywhere = the VPs are + // wrong (transposition or stale). + let vis18 = sun_vis_cascade(p0 + n0 * 0.02); + textureStore(out_hdr, px_full, vec4(vec3(vis18 * 50.0), 1.0)); + return; + } + if (debug == 19.0) { + // Numeric dump: cascade-0 shadow projection (ndc.xyz) + the + // stored atlas depth at the landing texel (-1 = out of range, + // -9 = degenerate w). Read back via pt_trace_dump.txt. + let pw19 = p0 + n0 * 0.02; + let clip19 = u.shadow_vps[0] * vec4(pw19, 1.0); + var out19 = vec4(-9.0); + if (abs(clip19.w) > 1e-6) { + let ndc19 = clip19.xyz / clip19.w; + let uv19 = vec2(ndc19.x * 0.5 + 0.5, 0.5 - ndc19.y * 0.5); + var stored19 = -1.0; + if (uv19.x >= 0.0 && uv19.x <= 1.0 && uv19.y >= 0.0 && uv19.y <= 1.0) { + let dims19 = vec2(textureDimensions(shadow_atlas_0)); + stored19 = textureLoad(shadow_atlas_0, vec2(uv19 * dims19), 0); + } + out19 = vec4(ndc19.xyz, stored19); + } + accum_out[gid.y * u.size.x + gid.x] = out19; + textureStore(out_hdr, px_full, vec4(0.1, 0.0, 0.2, 1.0)); + return; + } if (debug == 6.0 || debug == 7.0) { // PT-2 validation: trace the primary ray through the TLAS // (ignoring the G-buffer) and show interpolated attributes. @@ -847,12 +950,34 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { ign_at(px_full + vec2(17, 59), 0u), ); } + var view_cur = normalize(u.cam_pos.xyz - p0); var radiance = direct_light(p0 + n0 * 0.02, n0, albedo0 * (1.0 - metal_cur), sun_r2); + // Sun SPECULAR at the primary surface — the raster-style GGX + // highlight (D carries the 1/pi) masked by the same visibility. + // Without it every surface reads as dead-matte under the sun. + { + let ndl0 = max(dot(n0, u.sun_dir.xyz), 0.0); + if (ndl0 > 0.0) { + let vis0 = sun_visibility(p0 + n0 * 0.02, n0, sun_r2); + if (vis0 > 0.0) { + let hv = normalize(view_cur + u.sun_dir.xyz); + let ndv = max(dot(n0, view_cur), 1e-4); + let ndh = max(dot(n0, hv), 0.0); + let vdh = max(dot(view_cur, hv), 1e-4); + let alpha0 = max(rough_cur * rough_cur, 1e-3); + let a2 = alpha0 * alpha0; + let dd = ndh * ndh * (a2 - 1.0) + 1.0; + let dterm = a2 / (3.14159265 * dd * dd); + let f0s = mix(vec3(0.04), albedo0, metal_cur); + let spec = fresnel_schlick3(vdh, f0s) * dterm * v_smith(ndv, ndl0, alpha0); + radiance += spec * u.sun_color.rgb * ndl0 * vis0; + } + } + } var throughput = vec3(1.0); var origin = p0 + n0 * 0.02; var n_cur = n0; var alb_cur = albedo0; - var view_cur = normalize(u.cam_pos.xyz - p0); let max_bounces = u32(u.cfg.y); for (var b = 0u; b < max_bounces; b = b + 1u) { From 4384741e82868c3d5037d7b92d31120963e1f38d Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 09:41:05 +0200 Subject: [PATCH 15/16] feat(pt): decouple the realtime ray budget from the raster render scale The RT trace grid is now capped at ~0.5 Mpx (960x540) and the upsampler + owner-pixel mapping generalize from the hardcoded factor 2 to arbitrary trace-to-full ratios. Raising the raster render scale (0.5 -> 0.75 shipped in the shooter after re-measuring: 46 -> 30 fps raster, dramatically sharper at 4K) therefore sharpens the image without multiplying path-tracing cost. Progressive stays uncapped. --- native/shared/src/renderer/pt_pass.rs | 7 ++++++- native/shared/src/renderer/shaders/pt.rs | 14 +++++++++++--- 2 files changed, 17 insertions(+), 4 deletions(-) diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index 4ebef04..2f488aa 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -91,8 +91,13 @@ impl Renderer { // 2x2 sample phase rotates per frame so the temporal EMA // integrates full-res coverage over 4 frames. Progressive mode // stays full-res. + // The realtime trace grid is capped at ~0.5 Mpx (960x540) so + // raising the raster render scale sharpens the image without + // multiplying the ray budget — the upsampler handles arbitrary + // trace-to-full ratios. Progressive stays uncapped: quality is + // its entire point. let (trace_w, trace_h) = if self.pt_mode >= 2 { - (surf_w.div_ceil(2), surf_h.div_ceil(2)) + (surf_w.div_ceil(2).min(960), surf_h.div_ceil(2).min(540)) } else { (surf_w, surf_h) }; diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index 8701092..b144a39 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -613,8 +613,14 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { // have ext == size and phase 0, making this the identity. var px_full = px; if (u.ext.x > u.size.x) { + // Generalized ratio (the trace grid is budget-capped, so the + // factor is not necessarily 2): integer scale keeps each trace + // texel pinned to one owner pixel. px_full = min( - px * 2 + vec2(i32(u.ext.z), i32(u.ext.w)), + vec2( + px.x * i32(u.ext.x) / i32(u.size.x), + px.y * i32(u.ext.y) / i32(u.size.y), + ), vec2(i32(u.ext.x) - 1, i32(u.ext.y) - 1), ); } @@ -1310,8 +1316,10 @@ fn cs_final(@builtin(global_invocation_id) gid: vec3) { // geometry (whose taps all mismatch) softly averaged rather than // black. let zc = lin_depth_a(d); - let hx = clamp(px.x / 2, 0, hw - 1); - let hy = clamp(px.y / 2, 0, hh - 1); + // Generalized ratio mapping (trace grid is budget-capped, not + // always exactly half of full res). + let hx = clamp(px.x * hw / fw, 0, hw - 1); + let hy = clamp(px.y * hh / fh, 0, hh - 1); var sum = vec3(0.0); var wsum = 0.0; for (var dy = -1; dy <= 1; dy = dy + 1) { From 78c8f4d4bb10ff56c912d9c73f8b25bd8e68932d Mon Sep 17 00:00:00 2001 From: Ralph Kuepper Date: Tue, 14 Jul 2026 15:07:29 +0200 Subject: [PATCH 16/16] feat: PT-3b - canonical SVGF denoiser for the realtime path tracer Replaces the frozen-seed workaround stack (frozen RNG, spike clamp, despeckle, fixed sigma schedule) with the real architecture: rolling samples, temporal moments + history length in a new ping-pong buffer pair, bilinear per-tap-validated reprojection, surface-flip write-through keyed to the texel footprint depth window, five variance-guided a-trous iterations with squared-weight variance propagation, and first-iteration feedback into colour history. Two bugs the frozen seed had been masking, both fixed: - tlas_version bump (every node transform = every gameplay frame) reset realtime accumulation to 1 spp permanently; the reset is now progressive-only. - prev_vp uploaded transposed: mat4_multiply products are already in WGSL M*v layout (unlike mat4_invert outputs), so reprojection was garbage under any camera motion since PT-3 M1. New numeric debug views 20-23 (history heat, variance, SVGF state dump, reprojection dump). Verified against converged PROG ground truth; shimmer meter: RT 37.6/15.9pct vs raster baseline 49.5/16.7pct - the path traced mode is now temporally quieter than raster. See docs/pt/PT-3b-svgf-denoiser.md. Claude-Session: https://claude.ai/code/session_018574BdCfSjdpNK3WLgx1hP --- docs/pt/PT-3b-svgf-denoiser.md | 93 +++++ native/shared/src/renderer/mod.rs | 91 +++-- native/shared/src/renderer/pt_pass.rs | 214 ++++++----- native/shared/src/renderer/shaders/pt.rs | 469 ++++++++++++++++------- 4 files changed, 607 insertions(+), 260 deletions(-) create mode 100644 docs/pt/PT-3b-svgf-denoiser.md diff --git a/docs/pt/PT-3b-svgf-denoiser.md b/docs/pt/PT-3b-svgf-denoiser.md new file mode 100644 index 0000000..672cd11 --- /dev/null +++ b/docs/pt/PT-3b-svgf-denoiser.md @@ -0,0 +1,93 @@ +# PT-3b — SVGF denoiser for the realtime mode + +Status: **landed & verified** (2026-07-14, 760M / DX12+DXC). + +The user-facing complaint this fixes: realtime mode looked like "a dirty +lens with small ice crystals on it" — a static, screen-glued grain +pattern. The round-7 frozen sample sequence produced exactly that by +design (a wrong-but-stable 1-spp estimate, spatially filtered), and it +turned out to be masking two load-bearing bugs that had disabled ALL +temporal accumulation since PT-3 M1. + +## What the realtime mode is now (canonical SVGF, Schied et al. 2017) + +- **Rolling RNG in every mode.** The frozen sequence is gone; unbiased + fresh samples each frame are what the temporal estimator needs. +- **Temporal accumulation with moments.** New moments ping-pong buffers + (bindings 18/19) carry `(mu1, mu2, history length, raw depth)` per + trace texel. Color and moments accumulate with `alpha = + max(1/N, 0.1)` (cumulative average while young, EMA once mature; 0.1 + rather than the paper's 0.2 because the half-res 2-bounce sky lottery + is the dominant noise source and the deeper average halves the + residual mottle — direct sun is cascade-deterministic, so the slower + EMA only delays indirect changes). +- **Bilinear 2×2 reprojection with per-tap depth validation** (tight + relative tolerance, surface identity) replacing point sampling + the + old loose-tolerance workaround. +- **Surface-flip write-through.** At half res a trace texel's owner + pixel alternates between surfaces with TAA jitter (grass blade vs + ground). When no tap matches but the stored surface is still inside + the texel's current footprint depth window, the history is kept + VERBATIM and the sample dropped — one texel holds one surface's + history; the depth-guided upsampler routes texels to full-res pixels + by surface. Blending would leak blade lighting onto the ground + (gray-blue mottle); resetting would pin the texel at 1 spp (speckle). + Only a stored surface that left the footprint is a disocclusion. +- **Variance-guided à-trous** — five iterations (steps 1/2/4/8/16), B3 + 5×5 kernel, `sigma_l = 4·sqrt(gaussian3x3(variance))`, variance + filtered alongside with squared weights, depth-gradient edge stop. + Young history (< 4 frames) substitutes a spatial variance estimate + and floors it at 0.25 so disoccluded texels blur instead of surviving + as false edges. +- **First-iteration feedback:** iteration 1's output is copied back + over the accum buffer as next frame's color history (moments stay + raw) — the SVGF detail that makes the loop stable at 1 spp. +- Removed: history spike clamp, despeckle max-clamp, fixed sigma + schedule, frozen IGN. Kept (documented deviation): a single + irradiance-space firefly clamp at 4.0, as reference implementations + do. + +## The two masked bugs (both invisible under the frozen seed) + +1. **`tlas_version` reset nuked realtime history every frame.** + `pt_accum_count` was zeroed whenever the TLAS version changed — and + it changes on every node transform, i.e. every gameplay frame. + Correct scoping: full reset for progressive only; realtime's + per-texel depth validation already invalidates exactly what changed. + Found via the debug-20 history-length view (black everywhere). + +2. **`prev_vp` was uploaded transposed — reprojection was garbage under + any camera motion.** The engine has TWO matrix storage conventions: + `mat4_invert` outputs land transposed relative to WGSL `M*v` (hence + inv_vp's transpose-at-upload), while `mat4_multiply` products are + already in `M*v` layout (the shadow cascade VPs upload raw for the + same reason). Transposing the composed prev-VP collapsed every + reprojection into a ~40-texel band at screen centre with a + nonsensical depth. Found via the debug-23 numeric dump; after the + fix the dump shows sub-texel exact reprojection (pos tracks column + 1:1, depth pair agrees to 0.1%) with the title camera orbiting. + +## Verification + +- Debug views: 20 = history length heat, 21 = variance ×10, + 22 = numeric dump (n / variance / tap mass / depth), + 23 = numeric dump of reprojected position + depth pair. + All dump through `pt_trace_dump.txt` (BLOOM_PT_DEBUG env). +- Ground truth: converged PROG stills (25 s static) — smooth dark + ground under grass; RT now matches its structure (the earlier + gray-blue mottle was cross-surface leakage, not signal). +- Shimmer meter (two stills 1.5 s apart, grass band): + **RT 37.6 mean / 15.9 % sparkle vs raster 49.5 / 16.7 %** — the path + traced mode is temporally QUIETER than the raster baseline. +- Cost: ~15 fps RT vs ~17 before at 0.75 render scale (moments buffers + + two extra wavelet iterations + feedback copy). + +## Notes / open + +- Known hybrid-sun tradeoff: cascade shadows cannot resolve grass + micro-shadows, so RT ground under grass is brighter and flatter than + the traced-sun PROG reference. Deliberate (stability over contact + detail); traced-sun-in-RT would be a quality toggle, not a fix. +- The write-through's footprint window compares current-frame linear + depth against a stored prev-clip depth; under fast motion it degrades + toward disocclusion-reset, which is the safe direction. diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 5263ff2..b049a29 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -1278,6 +1278,12 @@ pub struct Renderer { /// WGSL. Ping-pong pair (PT-3 reprojection reads other pixels from /// the previous frame). Lazily (re)created at render extent. pt_accum_buffers: [Option; 2], + /// PT-3b SVGF — luminance moments + geometry side-channel, ping-pong + /// with the accum pair: (mu1, mu2, history length, raw depth) per + /// trace texel. The kernel validates reprojection taps and derives + /// the temporal variance from it; the à-trous passes read it for + /// depth edge-stopping and sky markers. + pt_moments_buffers: [Option; 2], /// Index of the buffer holding the PREVIOUS frame's result (the /// read side of this frame's dispatch). Flipped after dispatch. pt_accum_idx: usize, @@ -1309,21 +1315,22 @@ pub struct Renderer { /// Debug-16 numeric readback (see pt_pass.rs). pt_readback_buffer: Option, pt_dump_written: bool, - /// PT-3 — à-trous denoiser (realtime mode): two iterations over the - /// temporally-blended radiance, buffer→scratch then scratch→hdr. + /// PT-3b — SVGF wavelet filter (realtime mode): four à-trous + /// iterations (steps 1/2/4/8) on the trace grid, then a full-res + /// upsample+modulate pass. pt_atrous_mid_pipeline: Option, pt_atrous_final_pipeline: Option, pt_atrous_layout: Option, - /// Per-iteration uniform (step size + sigmas + extent). - pt_atrous_params_bufs: [wgpu::Buffer; 3], + /// Per-iteration uniform (step size + first-iteration flag + extent). + pt_atrous_params_bufs: [wgpu::Buffer; 6], pt_atrous_scratch: Option, pt_atrous_scratch2: Option, - /// Three iterations: bg1[i] = accum buffer i → scratch (step 1), - /// bg_mid2 = scratch → scratch2 (step 2), bg2 = scratch2 → hdr - /// (step 4, upsampling). Invalidated with pt_bg. - pt_atrous_bg1: [Option; 2], - pt_atrous_bg_mid2: Option, - pt_atrous_bg2: Option, + /// Bind groups indexed [written accum idx][stage 0..5]: stage 0 = + /// accum[written]→scratch, 1..4 ping-pong the scratches, 5 = final + /// upsample to hdr. All six bind moments[written] as the geometry + /// side-channel, hence the per-written-idx variants. Invalidated + /// with pt_bg. + pt_atrous_bgs: [[Option; 6]; 2], /// True when the PT kernel actually replaced this frame's scene /// colour. Mode 1 leaves the raster frame on screen until a few /// samples have accumulated (a moving camera = perpetual resets = @@ -5197,6 +5204,22 @@ impl Renderer { ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Comparison), count: None, }, + // PT-3b SVGF — moments ping-pong (18 = previous + // frame read side, 19 = this frame's output). + wgpu::BindGroupLayoutEntry { + binding: 18, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 19, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: false }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, ]; let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("pt_layout"), @@ -5302,6 +5325,16 @@ impl Renderer { multisampled: false, }, count: None, }, + // SVGF geometry side-channel (the kernel's + // moments buffer): depth edge-stopping, history + // length and sky markers. + wgpu::BindGroupLayoutEntry { + binding: 6, visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only: true }, + has_dynamic_offset: false, min_binding_size: None, + }, count: None, + }, ], }); let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { @@ -5325,26 +5358,22 @@ impl Renderer { } else { (None, None, None) }; - let pt_atrous_params_bufs = [ - device.create_buffer(&wgpu::BufferDescriptor { - label: Some("pt_atrous_params_0"), - size: 32, - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - mapped_at_creation: false, - }), - device.create_buffer(&wgpu::BufferDescriptor { - label: Some("pt_atrous_params_1"), - size: 32, - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - mapped_at_creation: false, - }), + // Six stages: five à-trous iterations + the final upsample. + let pt_atrous_params_bufs = std::array::from_fn::<_, 6, _>(|i| { device.create_buffer(&wgpu::BufferDescriptor { - label: Some("pt_atrous_params_2"), + label: Some(match i { + 0 => "pt_atrous_params_0", + 1 => "pt_atrous_params_1", + 2 => "pt_atrous_params_2", + 3 => "pt_atrous_params_3", + 4 => "pt_atrous_params_4", + _ => "pt_atrous_params_5", + }), size: 32, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, - }), - ]; + }) + }); let pt_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_uniform"), size: std::mem::size_of::() as u64, @@ -6960,6 +6989,7 @@ impl Renderer { pt_bg: [None, None], pt_uniform_buffer, pt_accum_buffers: [None, None], + pt_moments_buffers: [None, None], pt_accum_idx: 0, pt_accum_count: 0, pt_prev_vp: [[0.0; 4]; 4], @@ -6980,9 +7010,7 @@ impl Renderer { pt_atrous_params_bufs, pt_atrous_scratch: None, pt_atrous_scratch2: None, - pt_atrous_bg1: [None, None], - pt_atrous_bg_mid2: None, - pt_atrous_bg2: None, + pt_atrous_bgs: [[None, None, None, None, None, None], [None, None, None, None, None, None]], probe_trace_sdf_pipeline, probe_trace_sdf_layout, probe_trace_sdf_bg_cache: [None, None], @@ -7483,12 +7511,11 @@ impl Renderer { // the old extent. Same for the à-trous scratch + bgs. self.pt_bg = [None, None]; self.pt_accum_buffers = [None, None]; + self.pt_moments_buffers = [None, None]; self.pt_accum_count = 0; self.pt_atrous_scratch = None; self.pt_atrous_scratch2 = None; - self.pt_atrous_bg1 = [None, None]; - self.pt_atrous_bg_mid2 = None; - self.pt_atrous_bg2 = None; + self.pt_atrous_bgs = [[None, None, None, None, None, None], [None, None, None, None, None, None]]; self.hiz_linearize_bg_cache = None; self.occlusion.invalidate_bindings(); for slot in self.hiz_downsample_bg_cache.iter_mut() { diff --git a/native/shared/src/renderer/pt_pass.rs b/native/shared/src/renderer/pt_pass.rs index 2f488aa..fba551c 100644 --- a/native/shared/src/renderer/pt_pass.rs +++ b/native/shared/src/renderer/pt_pass.rs @@ -71,10 +71,18 @@ impl Renderer { let prev_vp_for_reproject = self.pt_prev_vp; self.pt_prev_vp = vp_unjittered; // Geometry changed under the accumulated image (door opened, - // enemy died) → history is a lie, restart. + // enemy died) → PROGRESSIVE history is a lie, restart. Realtime + // must NOT reset here: tlas_version bumps on every node + // transform — during gameplay that is every single frame, which + // silently pinned the SVGF history at 1 sample (found via the + // debug-20 history-length view; the frozen-seed era masked it). + // Mode 2's per-tap depth validation already rejects exactly the + // texels whose surface actually changed. if self.tlas_built_version != self.pt_last_tlas_version { self.pt_last_tlas_version = self.tlas_built_version; - self.pt_accum_count = 0; + if self.pt_mode == 1 { + self.pt_accum_count = 0; + } } // Progressive mode + camera in motion: the raster frame stays on // screen (kernel write threshold) and any sample traced now is @@ -131,10 +139,24 @@ impl Renderer { mapped_at_creation: false, })); } + // SVGF moments side-channel (mu1, mu2, history length, raw + // depth), ping-pong with the accum pair. wgpu zero-inits, + // and pt_accum_count = 0 marks the whole history invalid. + for (i, slot) in self.pt_moments_buffers.iter_mut().enumerate() { + *slot = Some(self.device.create_buffer(&wgpu::BufferDescriptor { + label: Some(if i == 0 { "pt_moments_a" } else { "pt_moments_b" }), + size: needed, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + } + // COPY_SRC: the first à-trous iteration's output is copied + // back over the accum buffer as next frame's colour history + // (SVGF feeds back the once-filtered signal). self.pt_atrous_scratch = Some(self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("pt_atrous_scratch"), size: needed, - usage: wgpu::BufferUsages::STORAGE, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false, })); self.pt_atrous_scratch2 = Some(self.device.create_buffer(&wgpu::BufferDescriptor { @@ -144,9 +166,7 @@ impl Renderer { mapped_at_creation: false, })); self.pt_bg = [None, None]; - self.pt_atrous_bg1 = [None, None]; - self.pt_atrous_bg_mid2 = None; - self.pt_atrous_bg2 = None; + self.pt_atrous_bgs = [[None, None, None, None, None, None], [None, None, None, None, None, None]]; self.pt_accum_count = 0; self.pt_accum_idx = 0; } @@ -185,14 +205,16 @@ impl Renderer { [m[0][2], m[1][2], m[2][2], m[3][2]], [m[0][3], m[1][3], m[2][3], m[3][3]], ]; - // Same transpose story for the reprojection matrix. - let p = &prev_vp_for_reproject; - let prev_vp_t = [ - [p[0][0], p[1][0], p[2][0], p[3][0]], - [p[0][1], p[1][1], p[2][1], p[3][1]], - [p[0][2], p[1][2], p[2][2], p[3][2]], - [p[0][3], p[1][3], p[2][3], p[3][3]], - ]; + // The reprojection VP uploads RAW — the opposite of inv_vp. The + // two matrix conventions coexist: mat4_invert outputs land + // transposed relative to WGSL's M*v (hence inv_vp's transpose + // above), while mat4_multiply products are already in M*v + // layout — the shadow cascade VPs upload raw for the same + // reason. Transposing this one collapsed every reprojection + // into a ~40-texel band at screen centre (debug-23 dump), so + // history never matched under camera motion — invisible in the + // frozen-seed era, which is why it survived since PT-3 M1. + let prev_vp_t = prev_vp_for_reproject; let params = PtParamsCpu { inv_vp: inv_vp_t, prev_vp: prev_vp_t, @@ -271,6 +293,10 @@ impl Renderer { wgpu::BindGroupEntry { binding: 15, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[1]) }, wgpu::BindGroupEntry { binding: 16, resource: wgpu::BindingResource::TextureView(&self.shadow_map.depth_views[2]) }, wgpu::BindGroupEntry { binding: 17, resource: wgpu::BindingResource::Sampler(&self.shadow_map.sampler) }, + // SVGF moments: read prev (paired with accum read + // side), write out (paired with the write side). + wgpu::BindGroupEntry { binding: 18, resource: self.pt_moments_buffers[i].as_ref().unwrap().as_entire_binding() }, + wgpu::BindGroupEntry { binding: 19, resource: self.pt_moments_buffers[1 - i].as_ref().unwrap().as_entire_binding() }, ]; self.pt_bg[i] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("pt_bg"), @@ -320,93 +346,101 @@ impl Renderer { let written_idx = 1 - self.pt_accum_idx; self.pt_accum_idx = written_idx; - // ---- PT-3: à-trous denoise (realtime mode only) ---- - // Two edge-aware iterations over the temporally-blended - // radiance: accum[written] → scratch (step 1), scratch → hdr - // (step 2). Progressive mode converges on its own and writes - // hdr directly from the kernel. + // ---- PT-3b: SVGF wavelet filter (realtime mode only) ---- + // Four variance-guided à-trous iterations on the trace grid + // (steps 1/2/4/8), then the full-res upsample+modulate pass. + // After iteration 1 the once-filtered signal is copied back + // over the accum buffer: SVGF feeds the first wavelet output + // into next frame's colour history (moments stay raw). This is + // what makes the temporal loop stable at 1 spp — raw history + // carries every spike forward, once-filtered history does not. + // Progressive mode converges on its own and writes hdr + // directly from the kernel. if self.pt_mode >= 2 && self.pt_debug == 0.0 && self.pt_atrous_mid_pipeline.is_some() && self.pt_atrous_scratch.is_some() { - // SVGF-style schedule: the first iteration barely edge-stops - // on luminance (1-2 spp variance IS the signal to remove); - // later iterations tighten so real shading detail survives. - // Depth stopping guards geometry edges throughout. - for (i, (step, sigma_luma)) in - [(1.0f32, 2.0f32), (2.0, 0.5), (4.0, 0.3)].iter().enumerate() - { + // p.y = 1.0 flags the FIRST iteration: it may substitute a + // spatial variance estimate where the history is young. + for (i, step) in [1.0f32, 2.0, 4.0, 8.0, 16.0, 1.0].iter().enumerate() { + let first = if i == 0 { 1.0f32 } else { 0.0 }; let p = [ - [*step, *sigma_luma, trace_w as f32, trace_h as f32], + [*step, first, trace_w as f32, trace_h as f32], [surf_w as f32, surf_h as f32, 0.0, 0.0], ]; self.queue.write_buffer(&self.pt_atrous_params_bufs[i], 0, bytemuck::bytes_of(&p)); } - if self.pt_atrous_bg1[written_idx].is_none() { - self.pt_atrous_bg1[written_idx] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("pt_atrous_bg1"), - layout: self.pt_atrous_layout.as_ref().unwrap(), - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[0].as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: self.pt_accum_buffers[written_idx].as_ref().unwrap().as_entire_binding() }, - wgpu::BindGroupEntry { binding: 2, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, - ], - })); + if self.pt_atrous_bgs[written_idx][0].is_none() { + let scratch = self.pt_atrous_scratch.as_ref().unwrap(); + let scratch2 = self.pt_atrous_scratch2.as_ref().unwrap(); + let accum_w = self.pt_accum_buffers[written_idx].as_ref().unwrap(); + let moments_w = self.pt_moments_buffers[written_idx].as_ref().unwrap(); + // Stage src → dst chain: accum→s1, then the scratches + // ping-pong; the final upsample reads the last-written + // scratch. cs_final never writes dst; it gets whichever + // scratch is not its src (RO+RW of one buffer in a + // single group fails validation). + let chain: [(&wgpu::Buffer, &wgpu::Buffer); 6] = [ + (accum_w, scratch), + (scratch, scratch2), + (scratch2, scratch), + (scratch, scratch2), + (scratch2, scratch), + (scratch, scratch2), + ]; + for (i, (src, dst)) in chain.iter().enumerate() { + self.pt_atrous_bgs[written_idx][i] = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("pt_atrous_bg"), + layout: self.pt_atrous_layout.as_ref().unwrap(), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[i].as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: src.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: dst.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, + wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, + wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, + wgpu::BindGroupEntry { binding: 6, resource: moments_w.as_entire_binding() }, + ], + })); + } } - if self.pt_atrous_bg_mid2.is_none() { - self.pt_atrous_bg_mid2 = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("pt_atrous_bg_mid2"), - layout: self.pt_atrous_layout.as_ref().unwrap(), - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[1].as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch.as_ref().unwrap().as_entire_binding() }, - wgpu::BindGroupEntry { binding: 2, resource: self.pt_atrous_scratch2.as_ref().unwrap().as_entire_binding() }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, - ], - })); + + { + let ts = profiler.compute_pass_timestamp_writes("pt_atrous"); + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("pt_atrous"), + timestamp_writes: ts, + }); + pass.set_pipeline(self.pt_atrous_mid_pipeline.as_ref().unwrap()); + pass.set_bind_group(0, self.pt_atrous_bgs[written_idx][0].as_ref().unwrap(), &[]); + pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); } - if self.pt_atrous_bg2.is_none() { - self.pt_atrous_bg2 = Some(self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("pt_atrous_bg2"), - layout: self.pt_atrous_layout.as_ref().unwrap(), - entries: &[ - wgpu::BindGroupEntry { binding: 0, resource: self.pt_atrous_params_bufs[2].as_entire_binding() }, - wgpu::BindGroupEntry { binding: 1, resource: self.pt_atrous_scratch2.as_ref().unwrap().as_entire_binding() }, - // dst is unused by cs_final; bind an accum - // buffer to satisfy the layout (cannot alias - // the scratch2 src binding — RO+RW of the same - // buffer in one group fails validation). - wgpu::BindGroupEntry { binding: 2, resource: self.pt_accum_buffers[0].as_ref().unwrap().as_entire_binding() }, - wgpu::BindGroupEntry { binding: 3, resource: wgpu::BindingResource::TextureView(&self.hdr_rt_view) }, - wgpu::BindGroupEntry { binding: 4, resource: wgpu::BindingResource::TextureView(&self.depth_view) }, - wgpu::BindGroupEntry { binding: 5, resource: wgpu::BindingResource::TextureView(&self.albedo_rt_view) }, - ], - })); + // History feedback: the pass split makes the copy legal + // (buffer copies cannot live inside a compute pass). + encoder.copy_buffer_to_buffer( + self.pt_atrous_scratch.as_ref().unwrap(), + 0, + self.pt_accum_buffers[written_idx].as_ref().unwrap(), + 0, + needed, + ); + { + let ts = profiler.compute_pass_timestamp_writes("pt_atrous2"); + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("pt_atrous2"), + timestamp_writes: ts, + }); + pass.set_pipeline(self.pt_atrous_mid_pipeline.as_ref().unwrap()); + for i in 1..5 { + pass.set_bind_group(0, self.pt_atrous_bgs[written_idx][i].as_ref().unwrap(), &[]); + pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); + } + pass.set_pipeline(self.pt_atrous_final_pipeline.as_ref().unwrap()); + pass.set_bind_group(0, self.pt_atrous_bgs[written_idx][5].as_ref().unwrap(), &[]); + pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); } - - let ts = profiler.compute_pass_timestamp_writes("pt_atrous"); - let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { - label: Some("pt_atrous"), - timestamp_writes: ts, - }); - // Mid passes filter on the trace grid (steps 1 and 2); the - // final pass runs at FULL resolution with step 4 (it - // upsamples when the grids differ). - pass.set_pipeline(self.pt_atrous_mid_pipeline.as_ref().unwrap()); - pass.set_bind_group(0, self.pt_atrous_bg1[written_idx].as_ref().unwrap(), &[]); - pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); - pass.set_bind_group(0, self.pt_atrous_bg_mid2.as_ref().unwrap(), &[]); - pass.dispatch_workgroups((trace_w + 7) / 8, (trace_h + 7) / 8, 1); - pass.set_pipeline(self.pt_atrous_final_pipeline.as_ref().unwrap()); - pass.set_bind_group(0, self.pt_atrous_bg2.as_ref().unwrap(), &[]); - pass.dispatch_workgroups((surf_w + 7) / 8, (surf_h + 7) / 8, 1); } // Mirrors the kernel's write threshold: mode 1 leaves the raster // frame on screen until 8 samples exist (u.size.w carried the @@ -419,7 +453,11 @@ impl Renderer { // Copies a window of the accum buffer (center of frame) into a // staging buffer each frame; the previous frame's copy is mapped // (blocking) and dumped to pt_trace_dump.txt once. - if (self.pt_debug == 16.0 || self.pt_debug == 17.0 || self.pt_debug == 19.0) + if (self.pt_debug == 16.0 + || self.pt_debug == 17.0 + || self.pt_debug == 19.0 + || self.pt_debug == 22.0 + || self.pt_debug == 23.0) && self.pt_accum_count > 30 && !self.pt_dump_written { diff --git a/native/shared/src/renderer/shaders/pt.rs b/native/shared/src/renderer/shaders/pt.rs index b144a39..4e30a44 100644 --- a/native/shared/src/renderer/shaders/pt.rs +++ b/native/shared/src/renderer/shaders/pt.rs @@ -97,9 +97,16 @@ struct InstanceGiData { // PT-3: ping-pong accumulation. Binding 8 = previous frame's buffer // (read), binding 13 = this frame's output. Reprojection reads OTHER // pixels from prev, which a single read_write buffer cannot do safely. +// +// Layout (SVGF): accum = (irradiance rgb, luminance variance); +// moments = (mu1, mu2, history length, raw depth). Progressive mode +// keeps its original (radiance sum, sample count) layout in accum and +// leaves the moments buffers untouched. @group(0) @binding(8) var accum: array>; @group(0) @binding(9) var out_hdr: texture_storage_2d; @group(0) @binding(13) var accum_out: array>; +@group(0) @binding(18) var moments: array>; +@group(0) @binding(19) var moments_out: array>; // Approximate linear view distance from the raw depth-buffer value // (GL-convention matrix, near 0.01: z_view ~= 2n / (1 - d)). Only used @@ -596,16 +603,13 @@ fn direct_light(p: vec3, n: vec3, alb: vec3, sun_r2: vec2) - fn cs_main(@builtin(global_invocation_id) gid: vec3) { if (gid.x >= u.size.x || gid.y >= u.size.y) { return; } let px = vec2(i32(gid.x), i32(gid.y)); - // Realtime mode uses a FROZEN random sequence (seed 0 instead of - // the frame index): every sample is deterministic per pixel, so - // the path tracer injects ZERO frame-to-frame noise — the à-trous - // smooths a static dither into a stable image. Progressive (and - // the debug views) keep the rolling sequence for convergence. - var seed_frame = u.size.z; - if (u.cfg.x >= 2.0 && u.cfg.w == 0.0) { - seed_frame = 0u; - } - rng_seed(gid.xy, seed_frame); + // A rolling per-frame sequence in EVERY mode: SVGF's temporal + // accumulation needs unbiased fresh samples each frame (a frozen + // sequence converges the EMA to a wrong-but-stable value that + // reads as static dirty-lens grain glued to the screen). The + // variance estimate below is what keeps rolling noise from + // shimmering: it tells the à-trous exactly where to blur hard. + rng_seed(gid.xy, u.size.z); // PT-3 half-res: realtime mode traces a half grid; map this trace // cell to its full-res G-buffer pixel (the 2x2 phase rotates per @@ -639,9 +643,12 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { if (is_sky(depth)) { // Leave the raster sky/clouds untouched. Realtime mode marks - // the texel (w = 1) so the a-trous passes skip it too. + // the texel as sky in the MOMENTS buffer (depth channel = far + // plane) so the a-trous passes and the upsampler skip it. if (u.cfg.x >= 2.0 && u.cfg.w == 0.0) { - accum_out[gid.y * u.size.x + gid.x] = vec4(0.0, 0.0, 0.0, 1.0); + let sky_idx = gid.y * u.size.x + gid.x; + accum_out[sky_idx] = vec4(0.0); + moments_out[sky_idx] = vec4(0.0, 0.0, 0.0, 1.0); } return; } @@ -942,18 +949,16 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { var metal_cur = mr0.r; var rough_cur = mr0.g; // Realtime mode samples the primary sun cone with structured IGN - // noise, FROZEN in time (frame 0): a per-frame rolling sequence - // makes every stochastically-shadowed pixel (grass, canopies) - // oscillate around its mean forever — the EMA never settles it and - // it reads as sparkle across the whole frame. A fixed per-pixel - // sample gives a static dither that the à-trous flattens into - // STABLE soft shadows. Progressive keeps rolling white noise for - // unbiased convergence. + // noise, rolling per frame: spatially well-distributed (a 5x5 + // filter averages it nearly flat, unlike white PCG noise) and + // temporally unbiased so the SVGF accumulation converges to the + // true mean. Under the hybrid cascade sun this path only matters + // when shadow maps are disabled. Progressive keeps white noise. var sun_r2 = rand_2f(); if (u.cfg.x >= 2.0) { sun_r2 = vec2( - ign_at(px_full, 0u), - ign_at(px_full + vec2(17, 59), 0u), + ign_at(px_full, u.size.z), + ign_at(px_full + vec2(17, 59), u.size.z), ); } var view_cur = normalize(u.cam_pos.xyz - p0); @@ -1056,25 +1061,10 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } } - // NaN/Inf guard + firefly cap so one bad sample cannot poison the - // accumulator forever. The cap scales with history: at 1 spp a - // single bright specular sample reads as a white dot on screen, so - // clamp hard; as the average deepens it can absorb real energy. + // NaN/Inf guard so one bad sample cannot poison the accumulator. if (radiance.r != radiance.r || radiance.g != radiance.g || radiance.b != radiance.b) { radiance = vec3(0.0); } - let luma = dot(radiance, vec3(0.2126, 0.7152, 0.0722)); - // Progressive: relax with accumulation depth (a deep average can - // absorb real energy). Realtime: FIXED — its EMA window stays ~7 - // frames no matter how large the frame count grows, so the cap - // must never relax or fireflies return as the count climbs. - var cap = 4.0 + f32(min(u.size.w, 28u)); - // Realtime cap is TIGHT (3): the à-trous luminance edge-stop - // preserves bright outliers as if they were edges, so any firefly - // that survives the cap survives the filter too. Scene luma sits - // around 0.5-1.5, so 3 leaves sunlit highlights alone. - if (u.cfg.x >= 2.0) { cap = 3.0; } - if (luma > cap) { radiance *= cap / luma; } // ---- accumulate --------------------------------------------------------- @@ -1085,66 +1075,197 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { var out: vec3; if (mode >= 2.0) { - // PT-3 M1: temporally REPROJECTED exponential history. The - // previous frame's result is fetched where THIS surface was - // last frame (world pos through prev_vp), so camera motion no - // longer smears screen-space history into fog. Disocclusions - // and depth mismatches reject history (alpha = 1); accepted - // history blends at 0.15. accum.w carries the raw depth the - // texel's surface wrote, validated against the reprojected - // depth of this frame's surface. - var alpha = 1.0; - var hist = vec3(0.0); + // SVGF temporal accumulation (Schied et al. 2017). History and + // output store IRRADIANCE (radiance demodulated by the primary + // albedo) so the wavelet passes filter lighting only; the + // final pass re-multiplies by the full-res G-buffer albedo. + var irr = radiance / max(albedo0, vec3(0.05)); + // Firefly clamp — the one practical deviation from the paper + // (reference implementations keep one too): it must bind in + // IRRADIANCE space, because dividing by a dark albedo + // amplifies radiance outliers up to 20x. Sunlit irradiance + // sits around 1-3; 4 leaves real highlights alone. + let irr_luma = dot(irr, vec3(0.2126, 0.7152, 0.0722)); + if (irr_luma > 4.0) { + irr *= 4.0 / irr_luma; + } + let l_new = min(irr_luma, 4.0); + + // Reprojection: 2x2 BILINEAR taps around the reprojected + // position, each tap validated for geometric consistency + // (relative linearized depth against the moments buffer). + // Point sampling here quantizes to whole trace texels and + // forced the old loose-tolerance workaround; weighted taps + // give sub-texel reprojection and a honest per-tap test. + var hist_rgb = vec3(0.0); + var hist_m1 = 0.0; + var hist_m2 = 0.0; + var hist_n = 0.0; + var wsum = 0.0; + // Footprint depth window (current frame): one trace texel + // covers ~3x3 full-res pixels; used below to tell a jitter + // surface-flip apart from a true disocclusion. + var fp_lo = 1e30; + var fp_hi = 0.0; + { + let rx = max(i32(u.ext.x) / i32(u.size.x), 1); + let ry = max(i32(u.ext.y) / i32(u.size.y), 1); + for (var sy = 0; sy <= 1; sy = sy + 1) { + for (var sx = 0; sx <= 1; sx = sx + 1) { + let sp = min( + px_full + vec2(sx * (rx - 1), sy * (ry - 1)), + vec2(i32(u.ext.x) - 1, i32(u.ext.y) - 1), + ); + let dz = depth_at(sp); + if (dz >= 0.9999999) { continue; } + let zl = lin_depth(dz); + fp_lo = min(fp_lo, zl); + fp_hi = max(fp_hi, zl); + } + } + } + var reproj_in_bounds = false; + var nearest_prev_idx = 0u; let clip_prev = u.prev_vp * vec4(p0, 1.0); if (u.size.w > 0u && clip_prev.w > 1e-4) { let ndc_prev = clip_prev.xyz / clip_prev.w; let uv_prev = vec2(ndc_prev.x * 0.5 + 0.5, 0.5 - ndc_prev.y * 0.5); if (uv_prev.x >= 0.0 && uv_prev.x < 1.0 && uv_prev.y >= 0.0 && uv_prev.y < 1.0) { - let ppx = vec2(uv_prev * vec2(f32(u.size.x), f32(u.size.y))); - let pidx = min(ppx.y, u.size.y - 1u) * u.size.x + min(ppx.x, u.size.x - 1u); - let ph = accum[pidx]; - let zl_hist = lin_depth(ph.w); + reproj_in_bounds = true; + let pos = uv_prev * vec2(f32(u.size.x), f32(u.size.y)) - 0.5; + let base = vec2(floor(pos)); + let fr = pos - floor(pos); + let np = vec2( + min(u32(max(base.x + i32(round(fr.x)), 0)), u.size.x - 1u), + min(u32(max(base.y + i32(round(fr.y)), 0)), u.size.y - 1u), + ); + nearest_prev_idx = np.y * u.size.x + np.x; let zl_here = lin_depth(ndc_prev.z); - // Tolerance is deliberately loose: at half-res the - // reprojected texel quantizes to 2 full pixels, and on - // depth-chaotic surfaces (grass) a tight test rejects - // almost every frame — pinning texels at 1 spp. The - // ghosting a loose test admits is invisible exactly - // where the depth is chaotic. - if (abs(zl_hist - zl_here) < 0.12 * max(zl_hist, zl_here) + 0.05) { - alpha = 0.1; - hist = ph.rgb; + if (debug == 23.0) { + // Reprojection dump: where this texel thinks it was + // last frame (trace-grid units) + the depth pair the + // acceptance test compares. Static camera => pos + // must equal the texel's own coordinates. + accum_out[idx] = vec4(pos.x, pos.y, zl_here, lin_depth(moments[nearest_prev_idx].w)); + moments_out[idx] = vec4(0.0, 0.0, 0.0, depth); + return; + } + // Tap test is TIGHT (surface identity). Cross-surface + // blending is the expensive error: it leaks bright + // blade-top lighting onto the ground below (gray-blue + // mottle). Surface flips are handled after the loop, + // not by widening this tolerance. + let tol = 0.1 * zl_here + 0.02; + for (var ty = 0; ty <= 1; ty = ty + 1) { + for (var tx = 0; tx <= 1; tx = tx + 1) { + let q = base + vec2(tx, ty); + if (q.x < 0 || q.y < 0 || q.x >= i32(u.size.x) || q.y >= i32(u.size.y)) { + continue; + } + let qidx = u32(q.y) * u.size.x + u32(q.x); + let m = moments[qidx]; + // Sky texels and depth-inconsistent taps carry + // another surface's lighting — skip them. + if (m.w >= 0.9999999) { + continue; + } + let zl_hist = lin_depth(m.w); + if (abs(zl_hist - zl_here) > tol) { + continue; + } + let wx = mix(1.0 - fr.x, fr.x, f32(tx)); + let wy = mix(1.0 - fr.y, fr.y, f32(ty)); + let wt = wx * wy + 1e-4; + hist_rgb += accum[qidx].rgb * wt; + hist_m1 += m.x * wt; + hist_m2 += m.y * wt; + hist_n += m.z * wt; + wsum += wt; + } } } } - // SVGF-style albedo demodulation: history and output store - // IRRADIANCE (radiance divided by the primary albedo) so the - // à-trous passes filter lighting only; the final pass - // re-multiplies by the FULL-RES G-buffer albedo, keeping - // texture detail crisp despite the half-res trace and the - // aggressive smoothing. - var irr = radiance / max(albedo0, vec3(0.05)); - // The firefly cap must bind IRRADIANCE, not radiance: dividing - // by a dark albedo amplifies up to 20x, so a radiance-space cap - // leaves outliers the filter sigmas can't absorb. Sunlit - // irradiance sits around 1-3; 4 leaves highlights alone. - let irr_luma = dot(irr, vec3(0.2126, 0.7152, 0.0722)); - if (irr_luma > 4.0) { - irr *= 4.0 / irr_luma; + + // No tap matched. One trace texel holds ONE surface's history; + // TAA jitter re-picks which surface the owner pixel sees each + // frame on sub-texel geometry (grass). If the STORED surface + // still exists in this texel's current footprint, this frame's + // sample simply belongs to the other surface: keep the history + // verbatim and drop the sample (blending would leak lighting + // across surfaces; resetting would pin the texel at 1 spp and + // fill the screen with speckle). The upsampler routes trace + // texels to full-res pixels by depth, so the other surface + // draws its lighting from neighbouring texels. Only a stored + // surface that has LEFT the footprint is a true disocclusion. + if (wsum <= 1e-3 && reproj_in_bounds) { + let mnp = moments[nearest_prev_idx]; + if (mnp.w < 0.9999999 && mnp.z > 0.0 && fp_hi > 0.0 && fp_lo < 1e29) { + let zl_st = lin_depth(mnp.w); + let wtol = 0.1 * zl_st + 0.02; + if (zl_st > fp_lo - wtol && zl_st < fp_hi + wtol) { + accum_out[idx] = accum[nearest_prev_idx]; + moments_out[idx] = mnp; + if (debug == 22.0) { + // Numeric dump: n / variance / 99 = flip path / depth. + accum_out[idx] = vec4(mnp.z, max(mnp.y - mnp.x * mnp.x, 0.0), 99.0, mnp.w); + } + if (debug == 20.0) { + textureStore(out_hdr, px_full, vec4(vec3(mnp.z / 32.0), 1.0)); + } + if (debug == 21.0) { + let v_st = max(mnp.y - mnp.x * mnp.x, 0.0); + textureStore(out_hdr, px_full, vec4(vec3(v_st * 10.0), 1.0)); + } + return; + } + } + } + + var n_hist = 0.0; + if (wsum > 1e-3) { + hist_rgb /= wsum; + hist_m1 /= wsum; + hist_m2 /= wsum; + n_hist = hist_n / wsum; } - // History-guided spike clamp: with valid history, one frame may - // not jump more than 50% + 0.25 above it — single-frame bright - // spikes ARE noise at 1-2 spp, and the à-trous edge-stop - // (center-relative) cannot remove an outlier CENTER. Darkening - // (a shadow appearing) stays unclamped so lighting changes - // propagate at full speed downward and EMA speed upward. - if (alpha < 1.0) { - irr = min(irr, hist * 1.5 + vec3(0.25)); + // Canonical blend: cumulative average while the history is + // young (alpha = 1/N), settling to EMA. The floor is 0.1 + // rather than the paper's 0.2: our trace is half-res with a + // 2-bounce sky lottery as the dominant noise source, and the + // deeper average halves the residual mottle. Direct sun comes + // from the raster cascades (deterministic), so the slower EMA + // only delays indirect/ambient changes (~10 frames). + let n_new = min(n_hist + 1.0, 32.0); + let alpha_c = max(1.0 / n_new, 0.1); + let out_irr = mix(hist_rgb, irr, alpha_c); + let m1 = mix(hist_m1, l_new, alpha_c); + let m2 = mix(hist_m2, l_new * l_new, alpha_c); + // Temporal luminance variance — the signal that drives the + // wavelet filter's luminance sigma. Young history makes this + // unreliable; the first à-trous iteration substitutes a + // spatial estimate when n < 4 (accum.w carries n via moments). + let variance = max(m2 - m1 * m1, 0.0); + accum_out[idx] = vec4(out_irr, variance); + moments_out[idx] = vec4(m1, m2, n_new, depth); + if (debug == 22.0) { + // Numeric dump: n / variance / accepted tap mass / depth. + accum_out[idx] = vec4(n_new, variance, wsum, depth); + } + if (debug == 20.0) { + // History length heat: white = full 32-frame history. + textureStore(out_hdr, px_full, vec4(vec3(n_new / 32.0), 1.0)); + } + if (debug == 21.0) { + // Variance view (x10 so typical values are visible). + textureStore(out_hdr, px_full, vec4(vec3(variance * 10.0), 1.0)); } - let out_irr = mix(hist, irr, alpha); - accum_out[idx] = vec4(out_irr, depth); return; } else { + // Progressive keeps its firefly cap, relaxing with + // accumulation depth (a deep average can absorb real energy). + let luma = dot(radiance, vec3(0.2126, 0.7152, 0.0722)); + let cap = 4.0 + f32(min(u.size.w, 28u)); + if (luma > cap) { radiance *= cap / luma; } // Progressive: plain running sum; count lives on the CPU. // Ping-pong read/write at the same index (static camera only). let sum = prev.rgb + radiance; @@ -1167,15 +1288,26 @@ fn cs_main(@builtin(global_invocation_id) gid: vec3) { } "#; -/// PT-3 — edge-aware à-trous denoiser for the realtime mode. Two -/// iterations (step 1 then 2) over the temporally-blended radiance: -/// `cs_mid` filters buffer→buffer, `cs_final` filters buffer→out_hdr. -/// Edge stopping = relative linearized depth (accum.w carries the raw -/// depth) + luminance similarity; sky texels (w = 1 marker) pass -/// through untouched and are never written to hdr. +/// PT-3b — SVGF wavelet filter (Schied et al. 2017) for the realtime +/// mode. Four `cs_mid` à-trous iterations (steps 1/2/4/8) run on the +/// trace grid over the temporally-accumulated irradiance; `cs_final` +/// joint-bilaterally upsamples to full resolution and re-modulates the +/// G-buffer albedo. Buffers: src/dst = (irradiance rgb, luminance +/// variance w); `geo` = the kernel's moments buffer (mu1, mu2, history +/// length, raw depth) — static across iterations, it carries the depth +/// for edge-stopping and the sky marker (depth = far plane). +/// +/// The luminance edge-stop is VARIANCE-DRIVEN: sigma_l scales with the +/// per-texel noise estimate, so grainy regions blur hard while +/// converged shading detail survives. Variance travels with the signal, +/// filtered by the squared weights, shrinking each iteration exactly as +/// the residual noise does. This replaces the old fixed sigma schedule, +/// the despeckle clamp and the history spike clamp — with a correct +/// variance estimate none of those are needed. pub(in crate::renderer) const PT_ATROUS_WGSL: &str = r#" struct AtrousParams { - // x = step (texels), y = sigma_luma, z = trace width, w = trace height + // x = step (texels), y = 1.0 on the FIRST iteration (enables the + // short-history spatial variance fallback), z/w = trace dims p: vec4, // x/y = full G-buffer dims (cs_final upsamples trace -> full when // they differ), z/w unused. @@ -1189,11 +1321,17 @@ struct AtrousParams { // Full-res G-buffer albedo: cs_final re-modulates the filtered // irradiance with it (SVGF demodulation keeps textures crisp). @group(0) @binding(5) var albedo_full: texture_2d; +// Kernel moments buffer: (mu1, mu2, history length, raw depth). +@group(0) @binding(6) var geo: array>; fn lin_depth_a(d: f32) -> f32 { return 0.02 / max(1.0 - d, 1e-6); } +fn luma_of(c: vec3) -> f32 { + return dot(c, vec3(0.2126, 0.7152, 0.0722)); +} + // B3-spline kernel weight for |offset| 0/1/2. fn kern(d: i32) -> f32 { let a = abs(d); @@ -1202,46 +1340,93 @@ fn kern(d: i32) -> f32 { return 0.0625; } -fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { +// SVGF luminance sigma (paper value). +const SIGMA_L: f32 = 4.0; + +fn filter_at(px: vec2, w: i32, h: i32, step: i32, first: bool) -> vec4 { let cidx = u32(px.y) * u32(w) + u32(px.x); - var center = src[cidx]; - if (center.w >= 0.9999999) { + let g_c = geo[cidx]; + let center = src[cidx]; + if (g_c.w >= 0.9999999) { return center; } - let zc = lin_depth_a(center.w); - var lc = dot(center.rgb, vec3(0.2126, 0.7152, 0.0722)); - - // Despeckle (first iteration only): the edge-stopping weights below - // are measured FROM the center, so an outlier center rejects its - // neighbors and survives every iteration. Clamp the center to its - // immediate neighborhood maximum before blending — isolated bright - // texels cannot exceed what any neighbor has seen. - if (step == 1) { - var nmax = 0.0; + let zc = lin_depth_a(g_c.w); + let lc = luma_of(center.rgb); + + // Center variance. Temporal variance from a young history (< 4 + // frames, e.g. right after a disocclusion) is meaningless — the + // paper substitutes a spatial luminance-variance estimate there. + var var_c = max(center.w, 0.0); + if (first && g_c.z < 4.0) { + var s1 = 0.0; + var s2 = 0.0; + var cnt = 0.0; for (var dy = -1; dy <= 1; dy = dy + 1) { for (var dx = -1; dx <= 1; dx = dx + 1) { - if (dx == 0 && dy == 0) { - continue; - } let q = px + vec2(dx, dy); if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { continue; } - let s = src[u32(q.y) * u32(w) + u32(q.x)]; - if (s.w >= 0.9999999) { + let qi = u32(q.y) * u32(w) + u32(q.x); + if (geo[qi].w >= 0.9999999) { continue; } - nmax = max(nmax, dot(s.rgb, vec3(0.2126, 0.7152, 0.0722))); + let lq = luma_of(src[qi].rgb); + s1 += lq; + s2 += lq * lq; + cnt += 1.0; } } - let lim = nmax * 1.5 + 0.05; - if (nmax > 0.0 && lc > lim) { - center = vec4(center.rgb * (lim / lc), center.w); - lc = lim; + if (cnt > 1.0) { + let mu = s1 / cnt; + var_c = max(var_c, max(s2 / cnt - mu * mu, 0.0)); + } + // Variance boost: a young history's estimate is unreliable in + // BOTH directions, and underestimating is the expensive error + // (a 1-spp outlier with variance 0 survives every iteration as + // a false edge). Floor it so fresh texels blend spatially until + // their temporal estimate matures. + var_c = max(var_c, 0.25); + } + + // 3x3 gaussian prefilter of the variance (paper 4.2): keeps a + // single hot texel from stopping its own smoothing. + var vsum = var_c * 0.25; + var vwsum = 0.25; + for (var dy = -1; dy <= 1; dy = dy + 1) { + for (var dx = -1; dx <= 1; dx = dx + 1) { + if (dx == 0 && dy == 0) { + continue; + } + let q = px + vec2(dx, dy); + if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { + continue; + } + let qi = u32(q.y) * u32(w) + u32(q.x); + if (geo[qi].w >= 0.9999999) { + continue; + } + let gw = select(0.0625, 0.125, dx == 0 || dy == 0); + vsum += max(src[qi].w, 0.0) * gw; + vwsum += gw; } } + let sigma_l_denom = SIGMA_L * sqrt(max(vsum / vwsum, 0.0)) + 1e-3; + + // Depth gradient (central differences on linear depth) scales the + // depth edge-stop so steep-slope surfaces stay connected while + // depth discontinuities still stop the filter (paper eq. 3). + var dzdx = 0.0; + var dzdy = 0.0; + if (px.x > 0 && px.x < w - 1) { + dzdx = (lin_depth_a(geo[cidx + 1u].w) - lin_depth_a(geo[cidx - 1u].w)) * 0.5; + } + if (px.y > 0 && px.y < h - 1) { + dzdy = (lin_depth_a(geo[cidx + u32(w)].w) - lin_depth_a(geo[cidx - u32(w)].w)) * 0.5; + } var sum = vec3(0.0); + var sum_v = 0.0; var wsum = 0.0; for (var dy = -2; dy <= 2; dy = dy + 1) { for (var dx = -2; dx <= 2; dx = dx + 1) { @@ -1249,26 +1434,30 @@ fn filter_at(px: vec2, w: i32, h: i32, step: i32) -> vec4 { if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { continue; } - var s = src[u32(q.y) * u32(w) + u32(q.x)]; - if (dx == 0 && dy == 0) { - s = center; // use the despeckled value for the center tap - } - if (s.w >= 0.9999999) { + let qi = u32(q.y) * u32(w) + u32(q.x); + let g_q = geo[qi]; + if (g_q.w >= 0.9999999) { continue; } - let zq = lin_depth_a(s.w); - let wz = exp(-abs(zq - zc) / (0.15 * max(zc, zq) + 0.02)); - let lq = dot(s.rgb, vec3(0.2126, 0.7152, 0.0722)); - let wl = exp(-abs(lq - lc) / max(ap.p.y, 1e-3)); + let s = src[qi]; + let zq = lin_depth_a(g_q.w); + let z_denom = abs(dzdx) * f32(abs(dx) * step) + + abs(dzdy) * f32(abs(dy) * step) + + 0.01 * zc + 1e-4; + let wz = exp(-abs(zq - zc) / z_denom); + let wl = exp(-abs(luma_of(s.rgb) - lc) / sigma_l_denom); let wgt = kern(dx) * kern(dy) * wz * wl; sum += s.rgb * wgt; + // Variance contracts with the SQUARED weights — the + // estimate shrinks exactly as the filtered noise does. + sum_v += max(s.w, 0.0) * wgt * wgt; wsum += wgt; } } if (wsum < 1e-6) { return center; } - return vec4(sum / wsum, center.w); + return vec4(sum / wsum, sum_v / (wsum * wsum)); } @compute @workgroup_size(8, 8, 1) @@ -1279,13 +1468,14 @@ fn cs_mid(@builtin(global_invocation_id) gid: vec3) { return; } let px = vec2(i32(gid.x), i32(gid.y)); - dst[gid.y * u32(w) + gid.x] = filter_at(px, w, h, i32(ap.p.x)); + dst[gid.y * u32(w) + gid.x] = filter_at(px, w, h, i32(ap.p.x), ap.p.y > 0.5); } -// Final pass runs at FULL resolution: filter-and-write when the trace -// grid is full-res, depth-guided joint-bilateral upsample when the -// trace grid is half-res. Sky pixels (full-res depth at far plane) are -// never written so the raster sky survives. +// Final pass runs at FULL resolution: depth-guided joint-bilateral +// upsample of the filtered irradiance, re-modulated by the full-res +// G-buffer albedo. Sky pixels (full-res depth at far plane) are never +// written so the raster sky survives. When the trace grid IS the full +// grid it degenerates to a plain modulate-and-write. @compute @workgroup_size(8, 8, 1) fn cs_final(@builtin(global_invocation_id) gid: vec3) { let fw = i32(ap.p2.x); @@ -1303,18 +1493,16 @@ fn cs_final(@builtin(global_invocation_id) gid: vec3) { let alb = max(textureLoad(albedo_full, px, 0).rgb, vec3(0.05)); if (hw == fw) { let cidx = gid.y * u32(fw) + gid.x; - if (src[cidx].w >= 0.9999999) { + if (geo[cidx].w >= 0.9999999) { return; } - let r = filter_at(px, fw, fh, i32(ap.p.x)); - textureStore(out_hdr_a, px, vec4(r.rgb * alb, 1.0)); + textureStore(out_hdr_a, px, vec4(src[cidx].rgb * alb, 1.0)); return; } - // Half-res upsample: 3x3 taps around the containing trace texel, - // weighted by kernel distance and relative linear-depth agreement - // with THIS full-res pixel. The epsilon keeps thin foreground - // geometry (whose taps all mismatch) softly averaged rather than - // black. + // 3x3 taps around the containing trace texel, weighted by kernel + // distance and relative linear-depth agreement with THIS full-res + // pixel. The epsilon keeps thin foreground geometry (whose taps + // all mismatch) softly averaged rather than black. let zc = lin_depth_a(d); // Generalized ratio mapping (trace grid is budget-capped, not // always exactly half of full res). @@ -1329,11 +1517,12 @@ fn cs_final(@builtin(global_invocation_id) gid: vec3) { if (qx < 0 || qy < 0 || qx >= hw || qy >= hh) { continue; } - let s = src[u32(qy) * u32(hw) + u32(qx)]; - if (s.w >= 0.9999999) { + let qi = u32(qy) * u32(hw) + u32(qx); + if (geo[qi].w >= 0.9999999) { continue; } - let wz = exp(-abs(lin_depth_a(s.w) - zc) / (0.08 * zc + 0.02)); + let s = src[qi]; + let wz = exp(-abs(lin_depth_a(geo[qi].w) - zc) / (0.08 * zc + 0.02)); let wgt = kern(dx) * kern(dy) * wz + 1e-5; sum += s.rgb * wgt; wsum += wgt;