Live: local lights guided by light tiles

docs/ROADMAP.md

@@ -2325,6 +2325,18 @@ Cornell box, Light Room and Bistro, means within 0.01 %. Live only the camera pa
 (as RTXPT; all-Sobol cost 2.1 ms of an 8.4 ms trace): cost and the live shots (bench-window) unchanged. The loop guard
 now skips counted loops. Next: blue-noise pixel distribution (Cycles 4.2), its own PR.
 
+**Local lights guided by light tiles (2026-11-15, after it).** The live kernel chose among shadow-casting local
+lights uniformly: with many of them, most behind walls, few candidates found a light that reaches the point. Now each
+camera pixel reports the light its shadow ray found, a pass makes each 8 x 8 tile's reports a list, and the next
+frame's paths take candidates from their tile's list beside the uniform ones (docs/RUNTIME.md "Light tiles"; the
+scheme MegaLights and RTXPT guide their draws by). Unbiased whatever a list holds. Measured on two new shots
+(`scenes/lamp-rooms.ts`, 96 lamps in 12 rooms; `lamp-rooms-wide.ts`, 768; bench-window `still-rooms`,
+`still-rooms-wide`), dE / flicker: 0.67 / 0.138 -> 0.63 / 0.069 and 1.01 / 0.126 -> 0.58 / 0.081, for 0.65 ms of
+trace and a 0.22 ms pass (of 9.2 ms a frame). (The wide shot's reference was rendered with a light tree, on another
+branch: main's offline chooses among the 768 lamps uniformly and is noisy at 4096 samples. Against main's own:
+dE 0.90 -> 0.71.) Scenes without more than 8 shadow-casting local lights compile and run as before (Bistro,
+Kabukichō: picture and cost unchanged).
+
 ## 2. Temporal fixes
 
 Targets (2026-11-08): bench-temporal, bench-ghosting and their UE MRQ references move to Bistro night (§ 0 step 7); the

docs/RUNTIME.md

@@ -1279,6 +1279,18 @@ offline, the frame live; 16 bits), ids are uncorrelated.
   shader against itself on the GPU.
 - The loop guard leaves counted loops (`for (var i = N; i < M; i++)`, literals) alone, as Tint does.
 
+## Light tiles
+
+Live, a scene with more than 8 shadow-casting local lights (`gpu::F_LISTS`) guides their resampling by what the last
+frame found (pathtrace.wgsl "Light tiles", live.wgsl `light_tiles`):
+- A camera hit's resampled shadow ray that reaches its light reports the light (the direct buffer's fourth
+  component).
+- `light_tiles`, the live passes' first, makes each 8 x 8 tile's reports a list: the lights reported, each once,
+  ascending, after the direct buffer's pixels (64 vec4s a tile). The tile grid shifts by a pixel offset each frame.
+- The next frame's paths take 2 candidates from their pixel's list at a vertex the camera sees and 1 past it, spread
+  over the list without repeats, beside the uniform candidates (8 and 3), which then stand for the lights the list
+  lacks. Each candidate is weighed by the lights it stands for: the estimate is the same sum whatever a list holds.
+
 ## What the kernel trusts
 
 The kernel follows indices it reads from buffers and loops over counts it is given; it checks none of them (a

renderer/shaders/live.wgsl

@@ -32,6 +32,7 @@ struct LiveParams {
     dof: vec4<f32>,    // depth of field: CoC radius of infinity (output px; 0 = off), focus distance (m), max radius (px), -
     jit: vec4<f32>,    // this frame's sub-pixel jitter (every traced sample's: x, y in 0 .. 1 of a traced pixel), -, -
     state: vec4<u32>,  // which half of the lasting buffers is this frame's (0 / 1; `this_half`), -, -, -
+    lt: vec4<u32>,     // the light tiles (`light_tiles`): this frame's grid offset x, y (traced pixels), the lists' first vec4 in `light_rw`, tiles a row
 };
 // color = diffuse lighting + first-hit distance, emit = directly seen emission / sky + roughness, spec = specular lighting
 // + reflection hit distance, normal.w = the media before the first hit (`medium_tr`, `medium_share`; see pathtrace.wgsl).
@@ -107,6 +108,9 @@ struct Pixel { color: vec4<f32>, albedo: vec4<f32>, normal: vec4<f32>, emit: vec
 // (`compose` adds it as it is, UE lights directly).
 @group(0) @binding(26) var<storage, read> direct_in: array<vec4<f32>>;
 @group(0) @binding(27) var<storage, read_write> lock_fwd: array<atomic<u32>>;
+// The kernel's direct buffer again, to write (`light_tiles` alone): a vec4 a traced pixel (w: the local light its
+// camera hit's resampled shadow ray found; pathtrace.wgsl "Light tiles"), then the tiles' lists.
+@group(0) @binding(28) var<storage, read_write> light_rw: array<vec4<f32>>;
 
 // Element i of this frame's half of a lasting buffer, and of the previous frame's.
 fn this_half(i: u32) -> u32 { return L.state.x * L.size.x * L.size.y + i; }
@@ -240,6 +244,49 @@ const RESOLVE_N: u32 = 8u;
 // w: the reflection's hit distance, weighted the same way (0: none), for the temporal pass's mirror-image
 // reprojection. A pixel's own ray is one random direction of the lobe (or none, when it took the diffuse lobe), so
 // its distance alone put the mirror image anywhere along the ray.
+// ---- The light tiles: each 8 x 8 tile's pixels' reports made its list (pathtrace.wgsl "Light tiles": the lights
+// reported, each once: index, reports; ascending; the 64th vec4 holds their number, 63 at most). A workgroup is a
+// tile: the grid is the traced one shifted by L.lt.xy, another offset each frame, so tile borders do not stay put.
+// (Float values, not integers' bits: pathtrace.wgsl LT_NONE says why.)
+const LT_NONE: u32 = 16777216u;
+var<workgroup> lt_keys: array<u32, 64>;
+var<workgroup> lt_first: array<u32, 64>;
+
+@compute @workgroup_size(8, 8, 1)
+fn light_tiles(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(local_invocation_index) li: u32, @builtin(workgroup_id) wg: vec3<u32>) {
+    let px = vec2<i32>(gid.xy) - vec2<i32>(L.lt.xy);
+    var key = LT_NONE;
+    if (px.x >= 0 && px.y >= 0 && px.x < i32(L.isz.x) && px.y < i32(L.isz.y)) {
+        key = min(u32(light_rw[u32(px.y) * L.isz.x + u32(px.x)].w), LT_NONE);
+    }
+    lt_keys[li] = key;
+    workgroupBarrier();
+    // Whether it is the first of the tile's reports of its light, and how many there are (four reports a step).
+    var same = vec4<u32>(0u);
+    var before = vec4<u32>(0u);
+    let k0 = vec4<u32>(key);
+    for (var j = 0u; j < 64u; j += 4u) {
+        let eq = vec4<u32>(lt_keys[j], lt_keys[j + 1u], lt_keys[j + 2u], lt_keys[j + 3u]) == k0;
+        same += vec4<u32>(eq);
+        before += vec4<u32>(eq & (vec4<u32>(j, j + 1u, j + 2u, j + 3u) < vec4<u32>(li)));
+    }
+    let first = key != LT_NONE && before.x + before.y + before.z + before.w == 0u;
+    lt_first[li] = u32(first);
+    workgroupBarrier();
+    // The first reports are the list: each after those of smaller keys.
+    var less = vec4<u32>(0u);
+    var all = vec4<u32>(0u);
+    for (var j = 0u; j < 64u; j += 4u) {
+        let f = vec4<u32>(lt_first[j], lt_first[j + 1u], lt_first[j + 2u], lt_first[j + 3u]);
+        less += f * vec4<u32>(vec4<u32>(lt_keys[j], lt_keys[j + 1u], lt_keys[j + 2u], lt_keys[j + 3u]) < k0);
+        all += f;
+    }
+    let list = L.lt.z + (wg.y * L.lt.w + wg.x) * 64u;
+    let at = less.x + less.y + less.z + less.w;
+    if (first && at < 63u) { light_rw[list + at] = vec4<f32>(f32(key), f32(same.x + same.y + same.z + same.w), 0.0, 0.0); }
+    if (li == 0u) { light_rw[list + 63u] = vec4<f32>(f32(min(all.x + all.y + all.z + all.w, 63u)), 0.0, 0.0, 0.0); }
+}
+
 @compute @workgroup_size(8, 8, 1)
 fn refl(@builtin(global_invocation_id) gid: vec3<u32>) {
     let iw = L.isz.x;

renderer/shaders/pathtrace.wgsl

@@ -78,6 +78,7 @@ const F_UNLIT: u32 = 1048576u; // MeshBasicMaterial (MAT_UNLIT)
 const F_BACKGROUND: u32 = 2097152u; // scene.background (`sky_mis`)
 const F_VCOLOR: u32 = 4194304u; // vertex colours (`vcolor`)
 const F_GRID: u32 = 33554432u; // density grids (`vgrid`)
+const F_LISTS: u32 = 134217728u; // light tiles (`lt_list_start`): a scene whose shadow-casting local lights are resampled
 fn has(f: u32) -> bool { return (FEAT & f) != 0u; }
 
 struct Params {
@@ -107,6 +108,8 @@ struct Params {
     procs: vec4<u32>,      // procedural geometry (F_PROC): its data's first word in env_cdf, object count, first record in `instances`, -
     aerial: vec4<f32>,     // the sky's aerial perspective (F_AERIAL): density multiple, Rayleigh scale, Mie scale, -
     background: vec4<f32>, // three.js scene.background (radiance), 1 (F_BACKGROUND)
+    ltile: vec4<u32>,      // live light tiles (`lt_list_start`): the lists' first vec4 in direct_out (0: none), tiles a row, the grid's offset when they were built (x | y << 16), candidates a vertex the camera sees takes from its list
+    ltile2: vec4<u32>,     // the traced size they were built at, -, -
 };
 
 struct TriAttr { n0: vec4<f32>, n1: vec4<f32>, n2: vec4<f32>, uv: vec4<f32>, t: vec4<f32>, x: vec4<f32> };
@@ -1684,6 +1687,49 @@ fn local_light(li: Light, x: vec3<f32>, u: vec2<f32>) -> LS {
 // ones too when there are few (one shadow ray each). Otherwise: resampled importance sampling (Talbot 2005) of RIS_M
 // uniform candidates by unshadowed contribution, then one shadow ray. `vol` = in-scattering in the fog (phase
 // function) instead of the surface BSDF. Surface results add their specular part to `dl_spec`.
+// Light tiles (live; what guides MegaLights' draws and RTXPT's): each camera pixel reports the local light its
+// resampled shadow ray found (direct_out's w: `lt_fb`, a light-table index; LT_NONE: none), live.wgsl `light_tiles`
+// makes each 8 x 8 tile's reports a list after the kernel's own data in direct_out (the lights found, each once:
+// index, reports; ascending; the last vec4 their number), and the next frame's paths take resampling candidates
+// from their pixel's tile's list as well, spread over it without repeats (all of it when it is short). The
+// candidates drawn uniformly, as before, then stand for the lights the list lacks (`lt_in_list`). A list is a
+// guide: whatever it holds the estimate is the same sum, so one gone stale (the table reordered, the camera moved
+// on) costs noise, not light. Compiled in for scenes whose shadow-casting local lights are resampled (F_LISTS).
+const LT_TILE: u32 = 8u;
+const LT_LIST: u32 = 64u;
+// (Indices and counts are stored as float values, not as their bits: a small integer's bits are a denormal float,
+// which a driver may flush to zero. So a light-table index stays below LT_NONE, 2^24.)
+const LT_NONE: u32 = 16777216u;
+var<private> lt_list: u32;
+var<private> lt_list_n: u32;
+var<private> lt_fb: u32;
+
+// The list of the tile pixel `px` was in when the lists were built, and how many lights it holds.
+fn lt_list_start(px: vec2<u32>) {
+    lt_list_n = 0u;
+    lt_fb = LT_NONE;
+    if (!LIVE || P.ltile.x == 0u) { return; }
+    let q = min(vec2<u32>(vec2<f32>(px) * vec2<f32>(P.ltile2.xy) / vec2<f32>(P.size.xy)), P.ltile2.xy - 1u);
+    let c = q + vec2<u32>(P.ltile.z & 0xffffu, P.ltile.z >> 16u);
+    lt_list = P.ltile.x + ((c.y / LT_TILE) * P.ltile.y + c.x / LT_TILE) * LT_LIST;
+    if (lt_list + LT_LIST > direct_len()) { return; }
+    lt_list_n = min(u32(direct_load(lt_list + LT_LIST - 1u).x), LT_LIST - 1u);
+}
+
+// Whether the pixel's list holds light `i`.
+fn lt_in_list(i: u32) -> bool {
+    var lo = 0u;
+    var hi = lt_list_n;
+    for (var k = 0u; k < 7u; k++) {
+        if (lo >= hi) { break; }
+        let mid = (lo + hi) / 2u;
+        let km = u32(direct_load(lt_list + mid).x);
+        if (km == i) { return true; }
+        if (km < i) { lo = mid + 1u; } else { hi = mid; }
+    }
+    return false;
+}
+
 fn local_lights(s: Surface, lb: Lobes, v: vec3<f32>, o: vec3<f32>, x: vec3<f32>, vol: bool, d: vec3<f32>, rd: u32) -> vec3<f32> {
     let nd = P.counts2.x;
     let nl = P.counts.x - nd;
@@ -1733,19 +1779,40 @@ fn local_lights(s: Surface, lb: Lobes, v: vec3<f32>, o: vec3<f32>, x: vec3<f32>,
     }
     // Target: unshadowed intensity x cosine (cheap; the full BSDF / phase is applied to the chosen light only).
     let m_cand = select(3u, RIS_M, seen_directly());
+    // With the pixel's list (light tiles), `nt` candidates more come from it, one in each `step` of its lights, and
+    // the uniform ones keep the lights it lacks. A uniform candidate stands for count / m_cand lights (weight 1 here,
+    // the factor applied at the end, as without a list), one of the list's for `step` (`listed`, relative to that).
+    var nt = 0u;
+    if (has(F_LISTS) && !vol && lt_list_n > 0u) { nt = min(select(P.ltile.w / 2u, P.ltile.w, seen_directly()), lt_list_n); }
+    let step = f32(lt_list_n) / f32(max(nt, 1u));
+    let listed = step * f32(m_cand) / f32(count);
+    var at = 0.0;
+    if (nt > 0u) { at = rng1(rng_sub(rd ^ 0x80000000u, 0x2000u)); }
     var wsum = 0.0;
     var pick = 0u;
     var pick_ph = 0.0;
     var pick_ls: LS;
-    for (var k = 0u; k < m_cand; k++) {
-        let r4 = rng4(rng_sub(rd ^ 0x80000000u, k));
-        let idx = first + min(u32(r4.x * f32(count)), count - 1u);
+    for (var k = 0u; k < m_cand + nt; k++) {
+        var r4: vec4<f32>;
+        var idx: u32;
+        var stands = 1.0;
+        if (k < m_cand) {
+            r4 = rng4(rng_sub(rd ^ 0x80000000u, k));
+            idx = first + min(u32(r4.x * f32(count)), count - 1u);
+        } else {
+            r4 = rng4(rng_sub(rd ^ 0x80000000u, 0x1000u + k));
+            idx = u32(direct_load(lt_list + min(u32((f32(k - m_cand) + at) * step), lt_list_n - 1u)).x);
+            if (idx < first || idx >= first + count) { continue; }
+            stands = listed;
+        }
         let ls = local_light(light_at(idx), x, r4.yz);
         var ph = luminance(ls.c);
         if (vol) { ph *= light_at(idx).x.x; } else { ph *= max(dot(s.n, ls.l), 0.0) + 0.05 * max(dot(s.ng, ls.l), 0.0); }
         if (ph <= 0.0) { continue; }
-        wsum += ph;
-        if (r4.w * wsum < ph) {
+        if (nt > 0u && k < m_cand && lt_in_list(idx)) { continue; }
+        let w = ph * stands;
+        wsum += w;
+        if (r4.w * wsum < w) {
             pick = idx;
             pick_ph = ph;
             pick_ls = ls;
@@ -1761,6 +1828,8 @@ fn local_lights(s: Surface, lb: Lobes, v: vec3<f32>, o: vec3<f32>, x: vec3<f32>,
                 let r = c * gtint(vis) * (wsum * f32(count) / (f32(m_cand) * pick_ph));
                 sum += r;
                 if (!vol) { dl_spec += r * split_frac; }
+                // The light tiles' report: the light the camera hit's ray found.
+                if (has(F_LISTS) && !vol && cur_depth == 0u) { lt_fb = pick; }
             }
         }
     }
@@ -3091,7 +3160,8 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let pix = gid.y * w + gid.x;
     // Every pixel's deterministic direct light is written each frame (0 unless the camera hits a surface): a pixel that
     // turned to sky must not keep what a surface left there.
-    if (LIVE && P.misc.w < 0.5 && direct_len() >= w * h) { direct_store(pix, vec4<f32>(0.0)); }
+    if (LIVE && P.misc.w < 0.5 && direct_len() >= w * h) { direct_store(pix, vec4<f32>(0.0, 0.0, 0.0, select(0.0, f32(LT_NONE), has(F_LISTS)))); }
+    if (has(F_LISTS)) { lt_list_start(gid.xy); }
     rng_start(gid.xy, P.size.z, P.size.w, 0u);
     cur_depth = 0u;
     let cam = rng4(rng_dim(D_CAMERA));
@@ -3394,7 +3464,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
         if (LIVE && depth == 0u && direct_len() >= P.size.x * P.size.y && !non_finite(dl_det) && !non_finite(dl_det_spec)) {
             radiance -= thr * dl_det;
             dl_spec = max(dl_spec - dl_det_spec, vec3<f32>(0.0));
-            if (P.misc.w < 0.5) { direct_store(gid.y * P.size.x + gid.x, vec4<f32>(thr * dl_det, 0.0)); }
+            if (P.misc.w < 0.5) { direct_store(gid.y * P.size.x + gid.x, vec4<f32>(thr * dl_det, select(0.0, f32(lt_fb), has(F_LISTS)))); }
         }
 
         let l = sample_bsdf(s, lb, v);

renderer/src/gpu.rs

@@ -115,6 +115,12 @@ pub struct Params {
     pub aerial: [f32; 4],
     /// three.js `scene.background` (F_BACKGROUND): its radiance, 1.
     pub background: [f32; 4],
+    /// The live light tiles (pathtrace.wgsl "Light tiles"; live/frame.rs): the lists' first vec4 in the direct
+    /// buffer (0: none this frame), tiles a row, the grid's offset when they were built (x | y << 16), candidates a
+    /// vertex the camera sees takes from its list beside the uniform ones (`live::LIST_CANDIDATES`; half past it).
+    pub ltile: [u32; 4],
+    /// The traced size the lists were built at, -, -.
+    pub ltile2: [u32; 4],
 }
 
 /// Room the live view's in-place geometry edits grow into (live/edits.rs `edit_geometry`), past the scene's own data:
@@ -1353,6 +1359,12 @@ pub const F_VCOLOR: u32 = 1 << 22;
 /// Density grids (`scene::GridBuild`): the grid atlas (binding 19) and `vgrid` for volume functions. (Bit 24 is
 /// F_AERIAL's.)
 pub const F_GRID: u32 = 1 << 25;
+/// More shadow-casting local lights than the kernel sums exactly at a camera hit (`RIS_EXACT`): they are resampled,
+/// guided live by the light tiles (pathtrace.wgsl "Light tiles"; the live view's `light_tiles` pass runs for such a
+/// scene).
+pub const F_LISTS: u32 = 1 << 27;
+/// Local lights the kernel sums exactly (pathtrace.wgsl `RIS_EXACT`; the test below holds the two together).
+pub const RIS_EXACT: usize = 8;
 /// The features that change the kernel's source (an `@if` flag) rather than switch code in it (`FEAT`).
 pub const F_SOURCE: u32 = F_UNITS;
 /// Software traversal's units (units.rs; shaders/trace_units.wgsl): the second level's walk, spliced into the
@@ -1388,7 +1400,7 @@ pub fn instance_features(scene: &Scene) -> u32 {
     // (F_GRID also for a volume function that calls `vgrid` on a node without a grid: it reads zeros.)
     let grid = !scene.grids.is_empty() || (!scene.media.is_empty() && scene.volume_fns.iter().any(|(_, code, helpers)| code.contains("vgrid(") || helpers.contains("vgrid(")));
     let media = if scene.media.is_empty() { 0 } else { F_MEDIA } | if scene.media.iter().any(|m| m.a[2] != 0) { F_HMEDIA } else { 0 } | if grid { F_GRID } else { 0 };
-    let particles = if scene.particles.iter().any(|p| !p.translucent) || !scene.gpu_particles.is_empty() { F_PEMIT } else { 0 } | if scene.procs.is_empty() { 0 } else { F_PROC } | if scene.tinted { F_TINT } else { 0 } | if scene.background.is_some() { F_BACKGROUND } else { 0 } | if scene.gpu_vcolors.is_empty() { 0 } else { F_VCOLOR };
+    let particles = if scene.particles.iter().any(|p| !p.translucent) || !scene.gpu_particles.is_empty() { F_PEMIT } else { 0 } | if scene.procs.is_empty() { 0 } else { F_PROC } | if scene.tinted { F_TINT } else { 0 } | if scene.background.is_some() { F_BACKGROUND } else { 0 } | if scene.gpu_vcolors.is_empty() { 0 } else { F_VCOLOR } | if scene.lights.iter().filter(|l| l.pos[3] != 0.0 && l.ext[3] >= 0.5).count() > RIS_EXACT { F_LISTS } else { 0 };
     if scene.gpu_instances.is_empty() {
         return media | particles;
     }
@@ -3361,14 +3373,15 @@ mod tests {
             ("F_MASKED", F_MASKED), ("F_FOG", F_FOG), ("F_VOLUME", F_VOLUME), ("F_TEX", F_TEX), ("F_FUNC", F_FUNC),
             ("F_GLASS", F_GLASS), ("F_COAT", F_COAT), ("F_HUGE", F_HUGE), ("F_INST", F_INST), ("F_MEDIA", F_MEDIA), ("F_PEMIT", F_PEMIT),
             ("F_HMEDIA", F_HMEDIA), ("F_PROC", F_PROC), ("F_TINT", F_TINT), ("F_UNLIT", F_UNLIT), ("F_BACKGROUND", F_BACKGROUND),
-            ("F_VCOLOR", F_VCOLOR), ("F_AERIAL", F_AERIAL), ("F_GRID", F_GRID),
+            ("F_VCOLOR", F_VCOLOR), ("F_AERIAL", F_AERIAL), ("F_GRID", F_GRID), ("F_LISTS", F_LISTS),
         ];
         let mut seen = 0u32;
         for (name, bit) in all {
             assert!(bit.is_power_of_two() && seen & bit == 0, "{name} = {bit:#x} is not a bit of its own");
             seen |= bit;
         }
         let wgsl = include_str!("../shaders/pathtrace.wgsl");
+        assert!(wgsl.contains(&format!("const RIS_EXACT: u32 = {RIS_EXACT}u;")), "RIS_EXACT: pathtrace.wgsl and gpu.rs differ");
         for line in wgsl.lines().filter(|l| l.starts_with("const F_")) {
             let name = &line[6..line.find(':').unwrap()];
             let value: u32 = line[line.find('=').unwrap() + 1..line.find("u;").unwrap()].trim().parse().unwrap();

renderer/src/lib.rs

@@ -743,6 +743,8 @@ pub fn prepare_as(hwrt: bool, units: bool, sc: &mut scene::Scene, timings: &mut
         lod_view: [0.0; 4],
         ltree: ltree_param,
         procs,
+        ltile: [0; 4],
+        ltile2: [0; 4],
     };
     let mut params = params;
     let mut first = 0u32;