Lighting on the GPU in the WebGPU path (t88 phase 2b)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

src/manifest.json

@@ -52,6 +52,7 @@
     "render/objects.js",
     "render/gpu/colour.js",
     "render/gpu/shaders.js",
+    "render/gpu/light.js",
     "render/gpu/device.js",
     "render/gpu/textures.js",
     "render/gpu/draw.js",

src/render/gpu/device.js

@@ -11,6 +11,9 @@ async function gpuInit(adapter) {
     GPU.context.configure({ device: d, format: GPU.format, alphaMode: 'opaque' });
     d.pushErrorScope('validation');
     GPU.colour = await d.createComputePipelineAsync({ layout: 'auto', compute: { module: d.createShaderModule({ code: GPU_COLOUR }), entryPoint: 'colour' } });
+    for (const [name, code, entryPoint] of [['lightRoof',GPU_LIGHT,'roof'],['lightSources',GPU_LIGHT,'sources'],['lightSpread',GPU_LIGHT_SPREAD,'main'],['lightPack',GPU_LIGHT_PACK,'main']]) {
+      GPU[name] = await d.createComputePipelineAsync({layout:'auto',compute:{module:d.createShaderModule({code}),entryPoint}});
+    }
     GPU.blur = await d.createComputePipelineAsync({ layout: 'auto', compute: { module: d.createShaderModule({ code: GPU_BLUR }), entryPoint: 'main' } });
     GPU.draw = await d.createRenderPipelineAsync({ layout: 'auto', vertex: { module: d.createShaderModule({ code: GPU_DRAW }), entryPoint: 'vs' }, fragment: { module: d.createShaderModule({ code: GPU_DRAW }), entryPoint: 'fs', targets: [{ format: GPU.format }] }, primitive: { topology: 'triangle-list' } });
     const error = await d.popErrorScope(); if (error) throw error;
@@ -18,9 +21,9 @@ async function gpuInit(adapter) {
     GPU.params = new Float32Array(32);
     GPU.sampler = d.createSampler({ magFilter: 'linear', minFilter: 'linear' });
     if (features.length) {
-      GPU.queries = d.createQuerySet({ type: 'timestamp', count: 2 });
+      GPU.queries = d.createQuerySet({ type: 'timestamp', count: 4 });
       GPU.resolve = d.createBuffer({ size: 256, usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC });
-      GPU.readback = d.createBuffer({ size: 16, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
+      GPU.readback = d.createBuffer({ size: 32, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
     }
     GPU.ok = true;
   } catch (e) { gpuFail(e.message); }

src/render/gpu/draw.js

@@ -30,7 +30,7 @@ function gpuOverrides() {
     for (let y = y0; y < y1; y++) for (let x = x0; x < x1; x++) cellPixel(x, y);
   }
 }
-function gpuDraw(lighting, glowOn) {
+function gpuDraw(lighting, glowOn, lightUpdate) {
   const upload = performance.now();
   gpuUpload(lighting, glowOn);
   if (GPU.canvas.width !== R.canvas.width || GPU.canvas.height !== R.canvas.height) {
@@ -45,18 +45,22 @@ function gpuDraw(lighting, glowOn) {
     const pass = encoder.beginComputePass(); pass.setPipeline(GPU.blur); pass.setBindGroup(0, GPU.bindBlur[i]);
     pass.dispatchWorkgroups(Math.ceil(R.glowC.width / 8), Math.ceil(R.glowC.height / 8)); pass.end();
   }
+  const lightTimed = lighting && (lightUpdate || GPU.lightFrame === undefined);
+  if (lightTimed) { gpuLight(encoder, timed); GPU.lightFrame = R.frame; }
   gpuPresent(encoder, timed);
-  if (timed) { encoder.resolveQuerySet(GPU.queries, 0, 2, GPU.resolve, 0); encoder.copyBufferToBuffer(GPU.resolve, 0, GPU.readback, 0, 16); }
+  if (timed) { encoder.resolveQuerySet(GPU.queries, 0, lightTimed ? 4 : 2, GPU.resolve, 0); encoder.copyBufferToBuffer(GPU.resolve, 0, GPU.readback, 0, lightTimed ? 32 : 16); }
   d.queue.submit([encoder.finish()]);
   if (timed) {
     GPU.timingBusy = true;
     GPU.timingPromise = GPU.readback.mapAsync(GPUMapMode.READ).then(() => {
       const ns = new BigUint64Array(GPU.readback.getMappedRange()); GPU.gpuMs = Number(ns[1] - ns[0]) / 1e6;
+      GPU.lightMs = lightTimed ? Number(ns[3] - ns[2]) / 1e6 : 0;
       GPU.readback.unmap(); GPU.timingBusy = false;
     }).catch(() => { GPU.timingBusy = false; });
   }
   if (PERF_MARKS || SET.showPerf) {
     perfMarks.override = perfMarks.color; perfMarks.upload = blit - upload; perfMarks.blit = performance.now() - blit;
+    perfMarks.gpuLight = GPU.lightMs || 0;
     if (GPU.gpuMs !== null) perfMarks.gpu = GPU.gpuMs;
   }
 }

src/render/gpu/light.js

@@ -0,0 +1,143 @@
+// GPU sources retain the CPU roof and transmission rules; directional relaxation
+// follows only left/up, then right/down, so it cannot introduce extra light paths.
+const GPU_LIGHT = `
+struct LightParams { grid: vec4f, world: vec4f, lamp: vec4f, levels: vec4f }
+@group(0) @binding(0) var<uniform> u: LightParams;
+@group(0) @binding(1) var<storage, read> lut: array<vec4f>;
+@group(0) @binding(2) var<storage, read> columns: array<vec4f>;
+@group(0) @binding(3) var matT: texture_2d<u32>;
+@group(0) @binding(4) var flagsT: texture_2d<u32>;
+@group(0) @binding(5) var partT: texture_2d<u32>;
+@group(0) @binding(6) var bgmT: texture_2d<u32>;
+@group(0) @binding(7) var machT: texture_2d<u32>;
+@group(0) @binding(8) var blockT: texture_2d<u32>;
+@group(0) @binding(9) var tempT: texture_2d<f32>;
+@group(0) @binding(10) var<storage, read_write> roofs: array<i32>;
+@group(0) @binding(11) var dest: texture_storage_2d<rgba16float, write>;
+fn phase(p: vec2i) -> i32 { return i32(lut[textureLoad(matT,p,0).r*3u+1u].x); }
+fn built(p: vec2i) -> bool {
+  return (phase(p)==${PH_SOLID} && (textureLoad(flagsT,p,0).r & ${F_NAT}u)==0u) ||
+    textureLoad(blockT,p/4,0).r!=0u || textureLoad(machT,p,0).r!=0u || textureLoad(partT,p,0).r!=0u;
+}
+fn loose(p: vec2i) -> bool {
+  for(var d=0;d<=12 && p.y+d<i32(u.world.y);d++) {
+    let q=p+vec2i(0,d); let ph=phase(q);
+    if(built(q)) { return true; }
+    if(ph!=${PH_POWDER} && ph!=${PH_LIQUID}) { return ph==${PH_EMPTY} || ph==${PH_GAS}; }
+  }
+  return false;
+}
+@compute @workgroup_size(64)
+fn roof(@builtin(global_invocation_id) id: vec3u) {
+  let lx=i32(id.x); if(lx>=i32(u.grid.z)) { return; }
+  let x=(i32(u.grid.x)+lx)*i32(u.world.z);
+  if(x<0 || x>=i32(u.world.x)) { roofs[lx]=0; return; }
+  let end=min(i32(u.world.y),i32((u.grid.y+u.grid.w)*u.world.z));
+  var y=i32(columns[x].y);
+  for(;y<end;y++) {
+    let p=vec2i(x,y); let ph=phase(p);
+    if((ph==${PH_SOLID} || ph==${PH_POWDER} || ph==${PH_LIQUID}) && !built(p) && !((ph==${PH_POWDER} || ph==${PH_LIQUID}) && loose(p))) { break; }
+  }
+  roofs[lx]=y;
+}
+@compute @workgroup_size(8,8)
+fn sources(@builtin(global_invocation_id) id: vec3u) {
+  let q=vec2i(id.xy); if(any(q>=vec2i(u.grid.zw))) { return; }
+  let p=(vec2i(u.grid.xy)+q)*i32(u.world.z); let res=i32(u.world.z);
+  var rgb=vec3f(0); var a=0.;
+  if(all(p>=vec2i(0)) && all(p<vec2i(u.world.xy))) {
+    let m=textureLoad(matT,p,0).r; let ph=phase(p);
+    var base=0.9;
+    if(built(p) || ((ph==${PH_POWDER} || ph==${PH_LIQUID}) && loose(p))) { base=0.978; }
+    else if(ph==${PH_EMPTY}) { base=select(0.978,0.96,textureLoad(bgmT,p,0).r!=0u); }
+    else if(ph==${PH_GAS}) { base=0.96; }
+    else if(ph==${PH_LIQUID}) { base=0.93; }
+    else if(m==${GLASS}u || m==${ICE}u) { base=0.96; }
+    a=pow(base,u.world.z);
+    if(p.y<roofs[q.x]) { rgb=u.levels.x*vec3f(1,0.98,0.94); }
+    else if(f32(p.y)<columns[p.x].x && ph!=${PH_SOLID}) { rgb=u.levels.x*0.9*vec3f(1,0.98,0.95); }
+    var hottest=0.;
+    for(var dy=0;dy<res && p.y+dy<i32(u.world.y);dy++) {
+      for(var dx=0;dx<res && p.x+dx<i32(u.world.x);dx++) { hottest=max(hottest,textureLoad(tempT,p+vec2i(dx,dy),0).r); }
+    }
+    if(hottest>520.) {
+      let gi=u32(clamp(i32((hottest-400.)/4.),0,699)); let glow=lut[${NMAT*3}u+gi];
+      rgb+=glow.rgb/255.*glow.a*1.35;
+    }
+  }
+  if(u.lamp.w>0.) {
+    let d=distance(vec2f(q),u.lamp.xy)/u.lamp.z;
+    if(d<1.) { rgb=max(rgb,(1.-d)*(1.-d)*1.05*vec3f(1,0.93,0.8)); }
+  }
+  textureStore(dest,q,vec4f(rgb,a));
+}`;
+const GPU_LIGHT_SPREAD = `
+@group(0) @binding(0) var source: texture_2d<f32>;
+@group(0) @binding(1) var dest: texture_storage_2d<rgba16float, write>;
+@group(0) @binding(2) var<uniform> direction: vec4i;
+@compute @workgroup_size(8,8)
+fn main(@builtin(global_invocation_id) id: vec3u) {
+  let p=vec2i(id.xy); let size=vec2i(textureDimensions(source)); if(any(p>=size)) { return; }
+  let own=textureLoad(source,p,0); var rgb=own.rgb;
+  let x=p+vec2i(direction.x,0); let y=p+vec2i(0,direction.x);
+  if(all(x>=vec2i(0)) && all(x<size)) { rgb=max(rgb,textureLoad(source,x,0).rgb*own.a); }
+  if(all(y>=vec2i(0)) && all(y<size)) { rgb=max(rgb,textureLoad(source,y,0).rgb*own.a); }
+  textureStore(dest,p,vec4f(rgb,own.a));
+}`;
+const GPU_LIGHT_PACK = `
+@group(0) @binding(0) var source: texture_2d<f32>;
+@group(0) @binding(1) var dest: texture_storage_2d<rgba8unorm, write>;
+@group(0) @binding(2) var<uniform> u: array<vec4f,4>;
+@compute @workgroup_size(8,8)
+fn main(@builtin(global_invocation_id) id: vec3u) {
+  let p=vec2i(id.xy); if(any(p>=vec2i(textureDimensions(source)))) { return; }
+  let rgb=floor(clamp(max(textureLoad(source,p,0).rgb,vec3f(u[3].y)),vec3f(0),vec3f(1))*255.)/255.;
+  textureStore(dest,p,vec4f(rgb,1));
+}`;
+function gpuLightGrid() {
+  const res = SET.lighting === 'fast' ? 4 : 2, M = Math.ceil(28 / res);
+  const lw = Math.ceil(R.bw / res) + 2 * M, lh = Math.ceil(R.bh / res) + 2 * M;
+  const key = `${lw}/${lh}/${res}`;
+  const update = GPU.lightKey !== key || GPU.lightWorld !== mat || GPU.lightFrame === undefined || R.frame % SET.lightEvery === 0;
+  if (update) {
+    GPU.lightKey = key; GPU.lightWorld = mat; GPU.lightFrame = R.frame;
+    GPU.lightGrid = [Math.floor(R.vx0 / res) - M, Math.floor(R.vy0 / res) - M, lw, lh, res];
+  }
+  [R.lx0, R.ly0, R.lw, R.lh, R.lres] = GPU.lightGrid;
+  return update;
+}
+function gpuLightBindings(group, view, buffer) {
+  const d = GPU.device;
+  if (GPU.lightUniform) { GPU.lightUniform.destroy(); GPU.lightRoofs.destroy(); }
+  GPU.lightUniform = d.createBuffer({ size: 64, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST });
+  GPU.lightRoofs = d.createBuffer({ size: Math.max(1,R.lw) * 4, usage: GPUBufferUsage.STORAGE });
+  const shared = [buffer(GPU.lightUniform), buffer(GPU.lut), buffer(GPU.columns),
+    ...[0,1,3,4].map(i => view(GPU.worldTextures[i])), view(GPU.mach), view(GPU.block), view(GPU.temp), buffer(GPU.lightRoofs)];
+  // Auto layouts contain only bindings used by each entry point.
+  GPU.bindRoof = d.createBindGroup({ layout: GPU.lightRoof.getBindGroupLayout(0), entries: shared.map((resource,binding)=>({binding,resource})).filter(e=>![6,9].includes(e.binding)) });
+  GPU.bindSources = group(GPU.lightSources, [...shared, view(GPU.lightWork[0])]);
+  if (!GPU.lightDirections) GPU.lightDirections = [-1,1].map(x => {
+    const b=d.createBuffer({size:16,usage:GPUBufferUsage.UNIFORM|GPUBufferUsage.COPY_DST});
+    d.queue.writeBuffer(b,0,new Int32Array([x,0,0,0])); return b;
+  });
+  GPU.bindSpread = GPU.lightDirections.map(b => [0,1].map(i => group(GPU.lightSpread, [view(GPU.lightWork[i]),view(GPU.lightWork[1-i]),buffer(b)])));
+  GPU.bindPack = group(GPU.lightPack, [view(GPU.lightWork[0]),view(GPU.lightTex),buffer(GPU.lightUniform)]);
+  GPU.lightFrame = undefined;
+}
+function gpuLight(encoder, timed) {
+  const q = GPU.device.queue, res = R.lres;
+  q.writeBuffer(GPU.lightUniform,0,new Float32Array([R.lx0,R.ly0,R.lw,R.lh,W,H,res,0,
+    (bot.x+bot.w/2)/res-R.lx0,(bot.y+bot.h/2)/res-R.ly0,34/res,+bot.lamp,
+    SET.daylight*(1-0.32*(R.rain||0)),SET.ambient,0,0]));
+  const pass = encoder.beginComputePass(timed ? {timestampWrites:{querySet:GPU.queries,beginningOfPassWriteIndex:2,endOfPassWriteIndex:3}} : {});
+  pass.setPipeline(GPU.lightRoof); pass.setBindGroup(0,GPU.bindRoof); pass.dispatchWorkgroups(Math.ceil(R.lw/64));
+  pass.setPipeline(GPU.lightSources); pass.setBindGroup(0,GPU.bindSources); pass.dispatchWorkgroups(Math.ceil(R.lw/8),Math.ceil(R.lh/8));
+  pass.setPipeline(GPU.lightSpread);
+  // 320 grains per direction: even maximum emission falls below one packed level.
+  const count = 320/res;
+  for(let direction=0;direction<2;direction++) for(let i=0;i<count;i++) {
+    pass.setBindGroup(0,GPU.bindSpread[direction][i&1]); pass.dispatchWorkgroups(Math.ceil(R.lw/8),Math.ceil(R.lh/8));
+  }
+  pass.setPipeline(GPU.lightPack); pass.setBindGroup(0,GPU.bindPack); pass.dispatchWorkgroups(Math.ceil(R.lw/8),Math.ceil(R.lh/8));
+  pass.end();
+}

src/render/gpu/textures.js

@@ -28,6 +28,10 @@ function gpuWorldTextures() {
   GPU.lut = gpuBuffer(tables);
   GPU.landData = new Float32Array((W + H) * 4);
   GPU.land = gpuBuffer(GPU.landData);
+  if (GPU.columns) GPU.columns.destroy();
+  GPU.columnData = new Float32Array(W * 4); GPU.columns = gpuBuffer(GPU.columnData);
+  GPU.mach = gpuTexture(W, H, 'r16uint'); GPU.block = gpuTexture(BW, BH, 'r8uint');
+  GPU.worldTextures.push(GPU.mach, GPU.block);
   GPU.bindColour = null;
 }
 function gpuViewTextures() {
@@ -40,15 +44,17 @@ function gpuViewTextures() {
   GPU.colourTex = gpuTexture(bw, bh, 'rgba8unorm', true);
   GPU.glowTex = gpuTexture(gw, gh, 'rgba8unorm', true);
   GPU.blurX = gpuTexture(gw, gh, 'rgba8unorm', true); GPU.blurY = gpuTexture(gw, gh, 'rgba8unorm', true);
-  GPU.lightTex = gpuTexture(Math.max(1, R.lw), Math.max(1, R.lh));
+  GPU.lightTex = gpuTexture(Math.max(1, R.lw), Math.max(1, R.lh), 'rgba8unorm', true);
-  GPU.viewTextures = [GPU.override, GPU.overrideGlow, GPU.colourTex, GPU.glowTex, GPU.blurX, GPU.blurY, GPU.lightTex];
+  GPU.lightWork = [0,1].map(() => gpuTexture(Math.max(1,R.lw),Math.max(1,R.lh),'rgba16float',true));
+  GPU.viewTextures = [GPU.override, GPU.overrideGlow, GPU.colourTex, GPU.glowTex, GPU.blurX, GPU.blurY, GPU.lightTex, ...GPU.lightWork];
   GPU.bindColour = null;
 }
 function gpuBindings() {
   if (GPU.bindColour) return;
   const d = GPU.device, buffer = b => ({ buffer: b }), view = t => t.createView();
   const group = (pipeline, resources) => d.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: resources.map((resource, binding) => ({ binding, resource })) });
-  GPU.bindColour = group(GPU.colour, [buffer(GPU.uniform), buffer(GPU.lut), buffer(GPU.land), ...GPU.worldTextures.map(view), view(GPU.override), view(GPU.overrideGlow), view(GPU.colourTex), view(GPU.glowTex)]);
+  GPU.bindColour = group(GPU.colour, [buffer(GPU.uniform), buffer(GPU.lut), buffer(GPU.land), ...GPU.worldTextures.slice(0,9).map(view), view(GPU.override), view(GPU.overrideGlow), view(GPU.colourTex), view(GPU.glowTex)]);
+  gpuLightBindings(group, view, buffer);
   GPU.bindDraw = group(GPU.draw, [buffer(GPU.uniform), view(GPU.colourTex), view(GPU.lightTex), view(GPU.blurY), GPU.sampler]);
   if (GPU.blurUniforms) for (const b of GPU.blurUniforms) b.destroy();
   GPU.blurUniforms = [0, 1].map(() => d.createBuffer({ size: 16, usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST }));
@@ -56,11 +62,21 @@ function gpuBindings() {
 }
 function gpuUpload(lighting, glowOn) {
   gpuWorldTextures(); gpuViewTextures(); gpuBindings();
-  const x = Math.max(0, R.vx0 - 2), y = Math.max(0, R.vy0 - 2);
+  const margin = lighting ? Math.ceil(28 / R.lres) * R.lres + R.lres : 2;
-  const w = Math.min(W, R.vx0 + R.bw + 2) - x, h = Math.min(H, R.vy0 + R.bh + 2) - y;
+  const x = Math.max(0, R.vx0 - margin), y = Math.max(0, R.vy0 - margin);
+  const w = Math.min(W, R.vx0 + R.bw + margin) - x, h = Math.min(H, R.vy0 + R.bh + margin + 12) - y;
   const arrays = [mat, flags, life, part, bgm, digp];
   arrays.forEach((a, i) => gpuWrite(GPU.worldTextures[i], a, W, x, y, w, h));
   gpuWrite(GPU.temp, temp, W, x, y, w, h);
+  if (lighting) {
+    let top=y;
+    for(let xx=x;xx<x+w;xx++) top=Math.min(top,roofY[xx]);
+    // Roof scans can start above the view; loose loads also inspect twelve grains below.
+    for(const [t,a] of [[GPU.worldTextures[0],mat],[GPU.worldTextures[1],flags],[GPU.worldTextures[3],part]]) gpuWrite(t,a,W,x,top,w,y+h-top);
+    gpuWrite(GPU.mach,mach,W,x,top,w,y+h-top);
+    const bx=x>>2, by=top>>2;
+    gpuWrite(GPU.block,blockP,BW,bx,by,Math.min(BW,Math.ceil((x+w)/4))-bx,Math.min(BH,Math.ceil((y+h)/4))-by);
+  }
   const fx = Math.max(0, Math.floor(x / FOG)), fy = Math.max(0, Math.floor(y / FOG));
   gpuWrite(GPU.worldTextures[7], fog, FW, fx, fy, Math.min(FW, Math.ceil((x + w) / FOG) + 1) - fx, Math.min(FH, Math.ceil((y + h) / FOG) + 1) - fy);
   // Small per-row/column tables may change in scene tools independently of world allocation.
@@ -71,7 +87,8 @@ function gpuUpload(lighting, glowOn) {
   q.writeBuffer(GPU.land, 0, land);
   gpuWrite(GPU.override, R.px, R.bw, 0, 0, R.bw, R.bh);
   gpuWrite(GPU.overrideGlow, R.glowPx, R.glowC.width, 0, 0, R.glowC.width, R.glowC.height);
-  if (lighting) gpuWrite(GPU.lightTex, R.lightPx, R.lw, 0, 0, R.lw, R.lh);
+  for(let i=0;i<W;i++) { GPU.columnData[i*4]=surfY[i]; GPU.columnData[i*4+1]=roofY[i]; }
+  q.writeBuffer(GPU.columns,0,GPU.columnData);
   const zp = R.zp, ox = Math.round(-(R.camX - R.vx0) * zp), oy = Math.round(-(R.camY - R.vy0) * zp);
   GPU.params.set([R.vx0, R.vy0, R.bw, R.bh, W, H, FW, FH,
     +SET.fog, R.rain, SET.grain, R.frame, +R.lens, Math.max(0, Object.keys(THERM).indexOf(SET.thermalPalette)), +SET.incandescence, +glowOn,

src/render/light.js

@@ -18,7 +18,7 @@ function computeLight() {
   const M = Math.ceil(28 / res);
   const lx0 = Math.floor(R.vx0 / res) - M, ly0 = Math.floor(R.vy0 / res) - M;
   const lw = Math.ceil(R.bw / res) + 2 * M, lh = Math.ceil(R.bh / res) + 2 * M;
-  if (lw !== R.lw || lh !== R.lh || res !== R.lres) {
+  if (lw !== R.lw || lh !== R.lh || res !== R.lres || !R.lr || R.lr.length !== lw * lh || R.lightC.width !== lw || R.lightC.height !== lh) {
     R.lw = lw; R.lh = lh; R.lres = res;
     R.lr = new Float32Array(lw * lh); R.lg = new Float32Array(lw * lh); R.lb = new Float32Array(lw * lh); R.latt = new Float32Array(lw * lh);
     R.lightC.width = lw; R.lightC.height = lh;

src/render/world.js

@@ -5,7 +5,9 @@ function fogColor(c, k) { return (c & 0xff000000) | (((c >>> 16 & 255) * k) << 1
 function drawWorld() {
   const perfStart = (PERF_MARKS || SET.showPerf) ? performance.now() : 0;
   if (PERF_MARKS || SET.showPerf) { perfMarks.light = 0; perfMarks.override = 0; delete perfMarks.gpu; }
+  const wasGpu = R.gpuOn;
   R.gpuOn = SET.renderer === 'webgpu' && SET.look === 'pixel' && GPU.ok;
+  if (R.gpuOn && !wasGpu) GPU.lightFrame = undefined;
   GPU.canvas.hidden = !R.gpuOn;
   if (R.gpuOn) R.ctx.clearRect(0, 0, R.canvas.width, R.canvas.height);
   const bw = R.bw, bh = R.bh;
@@ -33,7 +35,8 @@ function drawWorld() {
   }
   prepareSlopeBelts(vx0, vy0, vx0 + bw, vy0 + bh);
   const lighting = SET.lighting !== 'off' && !R.lens;
-  if (lighting && (R.frame % SET.lightEvery === 0 || !R.lw)) {
+  const lightUpdate = R.gpuOn && lighting ? gpuLightGrid() : false;
+  if (!R.gpuOn && lighting && (R.frame % SET.lightEvery === 0 || !R.lw || wasGpu)) {
     const lightStart = PERF_MARKS ? performance.now() : 0;
     computeLight();
     if (PERF_MARKS) perfMarks.light = performance.now() - lightStart;
@@ -48,7 +51,7 @@ function drawWorld() {
     perfMarks.color = performance.now() - colorStart;
     perfMarks.prepare = colorStart - perfStart - perfMarks.light;
   }
-  if (R.gpuOn) { gpuDraw(lighting, glowOn); return; }
+  if (R.gpuOn) { gpuDraw(lighting, glowOn, lightUpdate); return; }
   if (sm) { smoothDraw(lighting, glowOn); return; }
   const uploadStart = PERF_MARKS ? performance.now() : 0;
   R.bctx.putImageData(R.img, 0, 0);

src/ui/loop.js

@@ -101,7 +101,7 @@ function frame(now) {
   fpsSm = fpsSm * 0.95 + (1 / Math.max(0.001, dt)) * 0.05;
   if (SET.showPerf && R.frame % 15 === 0) {
     let aw = 0, ht = 0; for (let c = 0; c < NCH; c++) { if (awake[c]) aw++; if (hot[c]) ht++; }
-    $('perf').textContent = `${fpsSm.toFixed(0)} fps · frame ${frameMs.toFixed(1)} ms · sim ${simMs.toFixed(1)} ms/tick (move ${(perfTimes.simMove || 0).toFixed(1)}, heat ${(perfTimes.heatPass || 0).toFixed(1)}, rest ${Math.max(0, simMs - (perfTimes.simMove || 0) - (perfTimes.heatPass || 0)).toFixed(1)})\nworld ${(perfTimes.drawWorld || 0).toFixed(1)} ms · light ${(perfTimes.computeLight || 0).toFixed(1)} ms${R.gpuOn ? ` · overrides ${(perfMarks.override || 0).toFixed(1)} · upload ${(perfMarks.upload || 0).toFixed(1)} · GPU ${GPU.gpuMs === null ? "—" : GPU.gpuMs.toFixed(1)}` : ""} · overlays ${(perfTimes.overlays || 0).toFixed(1)} ms\nmoving areas ${aw} · hot areas ${ht} / ${NCH} · bars ${ents.length} · view ${R.bw}×${R.bh}`;
+    $('perf').textContent = `${fpsSm.toFixed(0)} fps · frame ${frameMs.toFixed(1)} ms · sim ${simMs.toFixed(1)} ms/tick (move ${(perfTimes.simMove || 0).toFixed(1)}, heat ${(perfTimes.heatPass || 0).toFixed(1)}, rest ${Math.max(0, simMs - (perfTimes.simMove || 0) - (perfTimes.heatPass || 0)).toFixed(1)})\nworld ${(perfTimes.drawWorld || 0).toFixed(1)} ms ${R.gpuOn ? "" : `· light ${(perfTimes.computeLight || 0).toFixed(1)} ms`}${R.gpuOn ? ` · overrides ${(perfMarks.override || 0).toFixed(1)} · upload ${(perfMarks.upload || 0).toFixed(1)} · GPU ${GPU.gpuMs === null ? "—" : GPU.gpuMs.toFixed(1)} ms (light ${(GPU.lightMs || 0).toFixed(1)})` : ""} · overlays ${(perfTimes.overlays || 0).toFixed(1)} ms\nmoving areas ${aw} · hot areas ${ht} / ${NCH} · bars ${ents.length} · view ${R.bw}×${R.bh}`;
   }
 }
 

test/t_gpu.py

@@ -72,6 +72,38 @@ with launch(viewport={'width': 960, 'height': 600}) as (browser, pg, url):
     assert pg.evaluate('!GPU.queries || GPU.gpuMs>=0')
     assert pg.evaluate('perfMarks.override>=0 && perfMarks.upload>=0')
     print('PASS separate override/upload clocks and optional GPU timestamps')
+    pg.evaluate('() => {' + (root / 'tools/scenes/t88-light.js').read_text() + '}')
+    # Draw GPU first with no CPU grid and forbid accidental CPU lighting calls.
+    pg.evaluate("""() => {
+      window.savedComputeLight=computeLight;
+      computeLight=()=>{throw Error('CPU light called in WebGPU');};
+      R.lr=R.lg=R.lb=R.latt=R.lightPx=null; R.lw=R.lh=0;
+      SET.renderer='webgpu';SET.look='pixel';SET.zoom=6;SET.lighting='full';SET.lightEvery=1;
+      resizeRender();bot.x=CAM[0]-bot.w/2;bot.y=CAM[1]-bot.h/2;snapCamera();drawWorld();
+      computeLight=savedComputeLight;
+    }""")
+    for mode in ['full', 'fast']:
+        for case in ['lamp', 'opening', 'magma', 'fire']:
+            parity = pg.evaluate("""([mode,kind]) => {
+              lightScene(kind);SET.lighting=mode;SET.lightEvery=1;R.frame=1000;R.rain=0;
+              SET.renderer='legacy';drawWorld();const a=worldPixelContext().getImageData(0,0,960,600).data;
+              R.rain=0;SET.renderer='webgpu';drawWorld();const b=worldPixelContext().getImageData(0,0,960,600).data;
+              let max=0,bad=0,sum=0;
+              for(let i=0;i<a.length;i+=4){let d=0;for(let j=0;j<3;j++){const e=Math.abs(a[i+j]-b[i+j]);sum+=e;d=Math.max(d,e);}max=Math.max(max,d);bad+=d>8;}
+              return {mode,kind,max,bad,mean:sum/(960*600*3)};
+            }""", [mode,case])
+            assert parity['bad'] == 0, parity
+            print('PASS light scene', parity)
+    cadence = pg.evaluate("""() => {
+      SET.lightEvery=3;R.frame=1002;drawWorld();const old=GPU.lightFrame,origin=R.lx0;
+      bot.lamp=!bot.lamp;R.camX+=4;R.frame=1003;drawWorld();
+      const held=GPU.lightFrame===old && R.lx0===origin;
+      R.frame=1005;drawWorld();const updated=GPU.lightFrame===1005 && R.lx0!==origin;
+      SET.lightEvery=1;return {held,updated};
+    }""")
+    assert all(cadence.values()), cadence
+    print('PASS GPU-only lighting, both resolutions and lightEvery cadence')
+    assert pg.evaluate('!GPU.queries || GPU.lightMs>=0')
     pg.evaluate('GPU.device.destroy()')
     pg.wait_for_function('!GPU.ok')
     assert pg.evaluate('SET.renderer==="legacy"')

tools/perf-measure.js

@@ -41,10 +41,13 @@ async function perfRender(lighting, overlays) {
       draw += perfLast.drawWorld / n;
       for (const [key, value] of Object.entries(perfMarks)) marks[key] = (marks[key] || 0) + value / n;
     }
+    let light = 0;
+    if (!R.gpuOn) {
     for (let k = 0; k < 5; k++) computeLight();
     const l0 = performance.now();
     for (let k = 0; k < n; k++) computeLight();
-    const light = (performance.now() - l0) / n;
+    light = (performance.now() - l0) / n;
+    }
     // Exercise the actual overlay calls and their arguments/conditions in frame().
     PERF_MARKS = false;
     for (let k = 0; k < 5; k++) frame(lastT + 1000 / 60);
@@ -58,7 +61,7 @@ async function perfRender(lighting, overlays) {
       await GPU.timingPromise;
       PERF_MARKS = true; drawWorld();
       await GPU.timingPromise;
-      PERF_MARKS = false; gpu_ms = GPU.gpuMs; marks.gpu = gpu_ms;
+      PERF_MARKS = false; gpu_ms = GPU.gpuMs; marks.gpu = gpu_ms; marks.gpuLight = GPU.lightMs;
     }
     return { frame_ms, drawWorld_ms: draw, computeLight_ms: light, marks_ms: marks,
       overlays: Object.fromEntries(Object.entries(data).map(([key, d]) => [key, { ms_per_frame: d.ms / n, calls: d.calls }])),

tools/scenes/t88-light.js

@@ -0,0 +1,25 @@
+// Frozen light comparisons, also used by the WebGPU browser regression suite.
+window.lightScene = kind => {
+  const cx = (W / 2) & ~3, cy = (H / 2) & ~3;
+  window.CAM = [cx, cy];
+  SET.fog = false; SET.shake = false; SET.cameraLead = false;
+  SET.daylight = kind === 'fire' ? 0 : 1; SET.ambient = 0.12;
+  bot.lamp = kind === 'lamp'; weather.next = 1e15; R.rain = 0;
+  const x0 = Math.max(0,cx-550), x1 = Math.min(W,cx+550);
+  const y0 = Math.max(0,cy-330), y1 = Math.min(H,cy+330);
+  for(let y=y0;y<y1;y++) for(let x=x0;x<x1;x++) {
+    const i=y*W+x;
+    const cave = x>cx-95 && x<cx+95 && y>cy-38 && y<cy+36;
+    const opening = kind==='opening' && x>=cx+24 && x<cx+40 && y<cy;
+    setCell(i, cave || opening || y<cy-100 ? EMPTY : STONE,20,F_NAT);
+    bgm[i]=part[i]=mach[i]=occ[i]=0; blockP[(y>>2)*BW+(x>>2)]=0;
+  }
+  for(let x=x0;x<x1;x++) { surfY[x]=cy-100; dirtY[x]=cy-90; graniteY[x]=cy+100; basaltY[x]=cy+200; }
+  if(kind==='magma' || kind==='fire') {
+    for(let y=cy+15;y<cy+23;y++) for(let x=cx-25;x<cx-8;x++) setCell(y*W+x,kind==='magma'?MAGMA:FIRE,kind==='magma'?1450:1100,F_NAT);
+  }
+  // A ledge makes occlusion/turning light paths visible around each source.
+  for(let y=cy-5;y<cy+7;y++) for(let x=cx+12;x<cx+45;x++) setCell(y*W+x,STONE,20,F_NAT);
+  updateRoof(); buildSky();
+};
+lightScene('lamp');