// Read float source channels, including transmittance, from the WGSL light pass.
function lightAt(a,x,y) { return a[(Math.floor(y/R.lres)-R.ly0)*R.lw+Math.floor(x/R.lres)-R.lx0]; }
async function testLight() {
  if (!W) return;
  renderDirty.fill(1);R.frame++;drawWorld();
  const d=GPU.device;
  if(!window.testLightPipe) window.testLightPipe=await d.createComputePipelineAsync({layout:'auto',compute:{module:d.createShaderModule({code:`
    @group(0) @binding(0) var source:texture_2d<f32>;
    @group(0) @binding(1) var<storage,read_write> out:array<vec4f>;
    @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id:vec3u){
      let size=textureDimensions(source);let k=id.x;if(k>=size.x*size.y){return;}
      out[k]=textureLoad(source,vec2i(i32(k%size.x),i32(k/size.x)),0);
    }`}),entryPoint:'main'}});
  const bytes=R.lw*R.lh*16, storage=d.createBuffer({size:bytes,usage:GPUBufferUsage.STORAGE|GPUBufferUsage.COPY_SRC});
  const read=d.createBuffer({size:bytes,usage:GPUBufferUsage.COPY_DST|GPUBufferUsage.MAP_READ});
  const bind=d.createBindGroup({layout:testLightPipe.getBindGroupLayout(0),entries:[{binding:0,resource:GPU.lightWork[0].createView()},{binding:1,resource:{buffer:storage}}]});
  const enc=d.createCommandEncoder(),pass=enc.beginComputePass();pass.setPipeline(testLightPipe);pass.setBindGroup(0,bind);pass.dispatchWorkgroups(Math.ceil(R.lw*R.lh/64));pass.end();
  enc.copyBufferToBuffer(storage,0,read,0,bytes);d.queue.submit([enc.finish()]);await read.mapAsync(GPUMapMode.READ);
  const data=new Float32Array(read.getMappedRange());
  [R.lr,R.lg,R.lb,R.latt]=[0,1,2,3].map(c=>Float32Array.from({length:R.lw*R.lh},(_,i)=>data[i*4+c]));
  read.unmap();read.destroy();storage.destroy();
}