diff --git a/docs/pt/PT-1-progressive-megakernel.md b/docs/pt/PT-1-progressive-megakernel.md new file mode 100644 index 0000000..daffd56 --- /dev/null +++ b/docs/pt/PT-1-progressive-megakernel.md @@ -0,0 +1,121 @@ +# 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. + +## 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 +bounce; no glass; sky at miss is the analytic Lumen sky, not the LUT sky. 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..8e1f0a2 --- /dev/null +++ b/docs/pt/PT-2-real-hit-shading.md @@ -0,0 +1,64 @@ +# 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. + +## M3 — GGX (landed b146b6d) + +`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 + 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/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/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/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/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/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/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 a59c502..b049a29 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,43 @@ 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)] +#[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. + 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/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 = 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], +} + #[repr(C)] #[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)] struct ProbeTemporalParams { @@ -586,6 +624,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)] @@ -1138,6 +1185,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. @@ -1208,6 +1262,83 @@ 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. + /// 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 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], + /// 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, + /// 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, + /// 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, + /// 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 + first-iteration flag + extent). + pt_atrous_params_bufs: [wgpu::Buffer; 6], + pt_atrous_scratch: Option, + pt_atrous_scratch2: 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 = + /// 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 /// dispatch time over the Hi-Z SW fallback when available. @@ -1817,6 +1948,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 { @@ -4908,6 +5048,346 @@ 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, 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 mut pt_layout_entries = vec![ + 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, + }, + // 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, + }, + // 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, + }, + // 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, + }, + // 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"), + 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: &pl_groups, + 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), tex_layout) + } 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, + }, + // 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, + }, + // 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 { + 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) + }; + // 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(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, + 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) @@ -6470,6 +6950,14 @@ impl Renderer { ssgi_intensity: 1.0, ssgi_radius: 20.0, ssgi_enabled: true, + // 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, @@ -6496,6 +6984,33 @@ impl Renderer { probe_trace_hw_pipeline, probe_trace_hw_layout, probe_trace_hw_bg_cache: [None, None], + pt_pipeline, + pt_layout, + 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], + 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, + 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_scratch2: 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], @@ -6991,6 +7506,16 @@ 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 BGs reference hdr_rt/depth views and the + // accumulation buffers are per-pixel-sized; both die with + // 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_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() { @@ -7235,6 +7760,37 @@ 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() + } + + /// 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`; @@ -7655,6 +8211,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, None]; resized = true; } @@ -7663,6 +8221,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; @@ -7670,6 +8237,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 @@ -7692,8 +8269,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, 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 = [ @@ -7702,6 +8333,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(), @@ -7758,6 +8404,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, None]; } // Populate TLAS instance slots. Clear stale entries from prior @@ -9571,6 +10219,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); @@ -9578,7 +10238,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..c6a531a 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_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 new file mode 100644 index 0000000..fba551c --- /dev/null +++ b/native/shared/src/renderer/pt_pass.rs @@ -0,0 +1,550 @@ +//! 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; + self.pt_wrote_frame = false; + 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_geo_vertex_buffer.is_none() + || 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 + // 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 + // 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; + } + // 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) → 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; + 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 + // 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; + } + + // ---- 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. + // 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).min(960), surf_h.div_ceil(2).min(540)) + } 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) ---- + // 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, + }; + if recreate { + 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, + })); + } + // 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 | wgpu::BufferUsages::COPY_SRC, + 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_bgs = [[None, None, None, None, None, None], [None, None, None, None, None, None]]; + self.pt_accum_count = 0; + self.pt_accum_idx = 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; + // 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]], + ]; + // 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, + 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: [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.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)); + + // ---- 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() }, + 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_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() }, + 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) }, + // 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"), + 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 ---- + { + 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[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((trace_w + 7) / 8, (trace_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; + + // ---- 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() + { + // 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, 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_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() }, + ], + })); + } + } + + { + 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); + } + // 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); + } + } + // 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 ---- + // 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 + || self.pt_debug == 22.0 + || self.pt_debug == 23.0) + && self.pt_accum_count > 30 + && !self.pt_dump_written + { + // 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"), + size: dump_pixels * 16, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + })); + encoder.copy_buffer_to_buffer( + // 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, + 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/mod.rs b/native/shared/src/renderer/shaders/mod.rs index 6092ed3..6e3d8d0 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, 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 new file mode 100644 index 0000000..4e30a44 --- /dev/null +++ b/native/shared/src/renderer/shaders/pt.rs @@ -0,0 +1,1536 @@ +//! 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) +//! 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 { + pos_range: vec4, // xyz world position, w = range + color_int: vec4, // rgb color, w = intensity +}; + +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 + 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 = 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, +}; + +// 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, + // 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; +@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; +// 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 +// 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. +// (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; +// 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; + +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) -------------------------------- + +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()); } + +// 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 +// trace grid may be half of that in realtime mode. +fn world_at(px: vec2, depth: f32) -> vec3 { + 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; + return w.xyz / w.w; +} + +fn depth_at(px: vec2) -> f32 { + let clamped = clamp(px, vec2(0), vec2(i32(u.ext.x) - 1, i32(u.ext.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); +} + +// ---- 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); + // 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); + 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); + 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) * 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; + } + // 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 / (1.0 - p_spec); + if (u.cfg.x >= 2.0) { + out.weight = min(out.weight, vec3(4.0)); + } + out.valid = true; + return out; +} + +// ---- 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). +// 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 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); + 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)); + // 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 + // 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) { + // 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( + 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), + ); + } + + let debug = u.cfg.w; + if (debug == 5.0) { + textureStore(out_hdr, px_full, vec4(1.0, 0.0, 1.0, 1.0)); + return; + } + + let depth = depth_at(px_full); + if (debug == 1.0) { + textureStore(out_hdr, px_full, vec4(vec3(depth), 1.0)); + return; + } + + if (is_sky(depth)) { + // Leave the raster sky/clouds untouched. Realtime mode marks + // 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) { + 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; + } + + 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_full, vec4(n0 * 0.5 + 0.5, 1.0)); + return; + } + if (debug == 3.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_full, 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_full, 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_full, 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_full, 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_full, 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_full, 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_full, 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_full, 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_out[idx16] = vec4( + h16.t, + f32(h16.instance_custom_data), + f32(h16.primitive_index), + f32(h16.kind), + ); + textureStore(out_hdr, px_full, 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_out[idx17] = vec4(p0, depth); + 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. + // 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_full, vec4(col, 1.0)); + return; + } + + // ---- one path sample -------------------------------------------------- + + // 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_full, 0).rg; + var metal_cur = mr0.r; + var rough_cur = mr0.g; + // Realtime mode samples the primary sun cone with structured IGN + // 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, u.size.z), + ign_at(px_full + vec2(17, 59), u.size.z), + ); + } + 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; + + let max_bounces = u32(u.cfg.y); + for (var b = 0u; b < max_bounces; b = b + 1u) { + 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)); + 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; + + // 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 { + 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). + 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), rand_2f()); + + 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) { + 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 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); + } + + // ---- 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) { + // 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) { + 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); + 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; + } + } + } + } + + // 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; + } + // 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)); + } + 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; + let n = f32(u.size.w) + 1.0; + 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 + // 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_full, vec4(out, 1.0)); +} +"#; + +/// 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 = 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. + 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; +// 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); + if (a == 0) { return 0.375; } + if (a == 1) { return 0.25; } + return 0.0625; +} + +// 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); + let g_c = geo[cidx]; + let center = src[cidx]; + if (g_c.w >= 0.9999999) { + return center; + } + 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) { + 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 lq = luma_of(src[qi].rgb); + s1 += lq; + s2 += lq * lq; + cnt += 1.0; + } + } + 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) { + let q = px + vec2(dx, dy) * step; + if (q.x < 0 || q.y < 0 || q.x >= w || q.y >= h) { + continue; + } + let qi = u32(q.y) * u32(w) + u32(q.x); + let g_q = geo[qi]; + if (g_q.w >= 0.9999999) { + continue; + } + 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, sum_v / (wsum * wsum)); +} + +@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), ap.p.y > 0.5); +} + +// 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); + 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 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 (geo[cidx].w >= 0.9999999) { + return; + } + textureStore(out_hdr_a, px, vec4(src[cidx].rgb * alb, 1.0)); + return; + } + // 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). + 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) { + 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 qi = u32(qy) * u32(hw) + u32(qx); + if (geo[qi].w >= 0.9999999) { + continue; + } + 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; + } + } + if (wsum < 1e-6) { + return; + } + textureStore(out_hdr_a, px, vec4((sum / wsum) * alb, 1.0)); +} +"#; 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/shared/src/renderer/ssgi_pass.rs b/native/shared/src/renderer/ssgi_pass.rs index 2e18295..62a2c40 100644 --- a/native/shared/src/renderer/ssgi_pass.rs +++ b/native/shared/src/renderer/ssgi_pass.rs @@ -27,7 +27,13 @@ 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). 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 568cdd1..5a09a8c 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_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 @@ -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_owns_frame() { let prev_idx = 1 - self.ssr_history_idx; let cur_idx = self.ssr_history_idx; 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(); 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);