lightingHEADmaster

7 files changed, 578 insertions, 26 deletions

diff --git a/src/content/materials/mod.rs b/src/content/materials/mod.rs
index 11c0cca..6338827 100644
--- a/src/content/materials/mod.rs
+++ b/src/content/materials/mod.rs
@@ -39,6 +39,10 @@ pub struct MaterialDef {
     pub density: u8,
     pub form: MaterialForm,
     pub sim_update: Option<fn(ctx: &mut UpdateCtx) -> PostUpdateAction>,
+
+    // lighting
+    pub emission: (u8, u8, u8, u8),
+    pub opacity: u8,
 }
 
 static MATERIALS: [MaterialDef; 11] = [
@@ -49,6 +53,9 @@ static MATERIALS: [MaterialDef; 11] = [
         // form usually but not always determines the update fn
         form: MaterialForm::Void,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0x00,
     },
     MaterialDef {
         name: "Sand",
@@ -56,6 +63,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 50,
         form: MaterialForm::Powder,
         sim_update: Some(powder::sim_update),
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xE0,
     },
     MaterialDef {
         name: "Dirt",
@@ -63,6 +73,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 55,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xE0,
     },
     MaterialDef {
         name: "Stone",
@@ -70,6 +83,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 115,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xFF,
     },
     MaterialDef {
         name: "Wood",
@@ -77,6 +93,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 60,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xFF,
     },
     MaterialDef {
         name: "Water",
@@ -84,6 +103,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 40,
         form: MaterialForm::Liquid,
         sim_update: Some(liquid::sim_update),
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0x70,
     },
     // TODO this is not a material
     MaterialDef {
@@ -93,6 +115,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 50,
         form: MaterialForm::Gas,
         sim_update: Some(fire::sim_update),
+
+        emission: (0xFF, 0x00, 0x00, 0x00),
+        opacity: 0x00,
     },
     MaterialDef {
         name: "Smoke",
@@ -100,6 +125,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 15,
         form: MaterialForm::Gas,
         sim_update: Some(gas::sim_update),
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0x10,
     },
     MaterialDef {
         name: "Steel",
@@ -107,6 +135,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 150,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xE0,
     },
     MaterialDef {
         name: "TNT",
@@ -114,6 +145,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 55,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xE0,
     },
     MaterialDef {
         name: "Cloth",
@@ -121,6 +155,9 @@ static MATERIALS: [MaterialDef; 11] = [
         density: 25,
         form: MaterialForm::Solid,
         sim_update: None,
+
+        emission: (0x00, 0x00, 0x00, 0x00),
+        opacity: 0xE0,
     },
 ];
 
diff --git a/src/proc_gen/herringbone.rs b/src/proc_gen/herringbone.rs
index a01c010..8d83750 100644
--- a/src/proc_gen/herringbone.rs
+++ b/src/proc_gen/herringbone.rs
@@ -3,7 +3,7 @@ use glam::IVec2;
 
 use crate::proc_gen::tileset_loader::TileOrientation;
 
-const SEED: u32 = 800;
+const SEED: u32 = 100000;
 
 pub fn split_position(pixel: IVec2, short: i32) -> (IVec2, IVec2, TileOrientation) {
     let grid = pixel.div_euclid(IVec2::splat(short));
diff --git a/src/renderer/instances.rs b/src/renderer/instances.rs
index 1292b99..7454bc7 100644
--- a/src/renderer/instances.rs
+++ b/src/renderer/instances.rs
@@ -38,18 +38,89 @@ struct RendererInstance {
     _padding: u32,
 }
 
+fn create_instance_bind_group(
+    device: &wgpu::Device,
+    layout: &wgpu::BindGroupLayout,
+    camera_uniform_buffer: &wgpu::Buffer,
+    palette_buffer: &wgpu::Buffer,
+    instance_buffer: &wgpu::Buffer,
+    cell_buffer: &wgpu::Buffer,
+    lighting_texture: &wgpu::TextureView,
+    sampler: &wgpu::Sampler,
+    lighting_uniform_buffer: &wgpu::Buffer,
+) -> wgpu::BindGroup {
+    device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: Some("Instance bind group"),
+        layout,
+        entries: &[
+            wgpu::BindGroupEntry {
+                binding: 0,
+                resource: camera_uniform_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 1,
+                resource: palette_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 2,
+                resource: instance_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 3,
+                resource: cell_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 4,
+                resource: wgpu::BindingResource::TextureView(lighting_texture),
+            },
+            wgpu::BindGroupEntry {
+                binding: 5,
+                resource: wgpu::BindingResource::Sampler(sampler),
+            },
+            wgpu::BindGroupEntry {
+                binding: 6,
+                resource: lighting_uniform_buffer.as_entire_binding(),
+            },
+        ],
+    })
+}
+
 pub struct InstancesRendererState {
     instance_pipeline: wgpu::RenderPipeline,
+    instance_bind_group_layout: wgpu::BindGroupLayout,
+    camera_uniform_buffer: wgpu::Buffer,
+    palette_buffer: wgpu::Buffer,
     instance_buffer: wgpu::Buffer,
     cell_buffer: wgpu::Buffer,
     instance_bind_group: wgpu::BindGroup,
     entity_capacity: usize,
+    sampler: wgpu::Sampler,
 
     renderer_entities: FxHashMap<EntityId, usize>,
     instances_count: usize,
 }
 
 impl InstancesRendererState {
+    /// Rebind the lighting resources, needed whenever the lighting texture is recreated.
+    pub fn rebuild_bind_group(
+        &mut self,
+        device: &wgpu::Device,
+        lighting_texture: &wgpu::TextureView,
+        lighting_uniform_buffer: &wgpu::Buffer,
+    ) {
+        self.instance_bind_group = create_instance_bind_group(
+            device,
+            &self.instance_bind_group_layout,
+            &self.camera_uniform_buffer,
+            &self.palette_buffer,
+            &self.instance_buffer,
+            &self.cell_buffer,
+            lighting_texture,
+            &self.sampler,
+            lighting_uniform_buffer,
+        );
+    }
+
     pub fn update_buffers(
         &mut self,
         device: &wgpu::Device,
@@ -185,6 +256,8 @@ impl InstancesRendererState {
         surface_config: &wgpu::SurfaceConfiguration,
         camera_uniform_buffer: &wgpu::Buffer,
         palette_buffer: &wgpu::Buffer,
+        lighting_texture: &wgpu::TextureView,
+        lighting_uniform_buffer: &wgpu::Buffer,
     ) -> Self {
         let instance_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("Instance shader"),
@@ -239,6 +312,35 @@ impl InstancesRendererState {
                         count: None,
                         visibility: wgpu::ShaderStages::FRAGMENT,
                     },
+                    // lighting
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::Texture {
+                            multisampled: false,
+                            sample_type: wgpu::TextureSampleType::Float { filterable: true },
+                            view_dimension: wgpu::TextureViewDimension::D2Array,
+                        },
+                        binding: 4,
+                        count: None,
+                        visibility: wgpu::ShaderStages::FRAGMENT,
+                    },
+                    // sampler
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
+                        binding: 5,
+                        count: None,
+                        visibility: wgpu::ShaderStages::FRAGMENT,
+                    },
+                    // lighting uniform
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Uniform,
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        binding: 6,
+                        count: None,
+                        visibility: wgpu::ShaderStages::FRAGMENT,
+                    },
                 ],
             });
 
@@ -288,35 +390,31 @@ impl InstancesRendererState {
 
         let cell_buffer = create_cell_buffer(&device, entity_capacity + CHUNK_SLOTS);
 
-        let instance_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
-            label: None,
-            layout: &instance_bind_group_layout,
-            entries: &[
-                wgpu::BindGroupEntry {
-                    binding: 0,
-                    resource: camera_uniform_buffer.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 1,
-                    resource: palette_buffer.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 2,
-                    resource: instance_buffer.as_entire_binding(),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 3,
-                    resource: cell_buffer.as_entire_binding(),
-                },
-            ],
+        let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
+            ..Default::default()
         });
+        let instance_bind_group = create_instance_bind_group(
+            device,
+            &instance_bind_group_layout,
+            camera_uniform_buffer,
+            palette_buffer,
+            &instance_buffer,
+            &cell_buffer,
+            lighting_texture,
+            &sampler,
+            lighting_uniform_buffer,
+        );
 
         InstancesRendererState {
             instance_pipeline,
+            instance_bind_group_layout,
+            camera_uniform_buffer: camera_uniform_buffer.clone(),
+            palette_buffer: palette_buffer.clone(),
             instance_buffer,
             cell_buffer,
             instance_bind_group,
             entity_capacity,
+            sampler,
 
             renderer_entities: FxHashMap::default(),
             instances_count: 0,
diff --git a/src/renderer/lighting.rs b/src/renderer/lighting.rs
new file mode 100644
index 0000000..94c7fb2
--- /dev/null
+++ b/src/renderer/lighting.rs
@@ -0,0 +1,333 @@
+use glam::{IVec2, Vec2};
+use wgpu::TextureViewDescriptor;
+
+use crate::{
+    camera::Camera,
+    content::materials::{MaterialDef, MaterialId},
+    sim::sim_manager::SimManager,
+};
+
+// light resolution = light pixels per cell
+const LIGHT_RESOLUTION: f32 = 1.0;
+const LAYERS: u32 = 4;
+const WORKGROUP_SIZE: u32 = 8;
+
+#[repr(C)]
+#[derive(Copy, Clone, bytemuck::Pod, bytemuck::Zeroable)]
+struct LightingUniform {
+    // world position of the lighting texture's (0, 0) texel
+    origin: [f32; 2],
+    _padding: [f32; 2],
+}
+
+fn create_scene_texture(device: &wgpu::Device, size: IVec2) -> wgpu::Texture {
+    device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("Scene texture"),
+        size: wgpu::Extent3d {
+            width: size.x as u32,
+            height: size.y as u32,
+            depth_or_array_layers: 1,
+        },
+        usage: wgpu::TextureUsages::STORAGE_BINDING
+            | wgpu::TextureUsages::TEXTURE_BINDING
+            | wgpu::TextureUsages::COPY_DST,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        dimension: wgpu::TextureDimension::D2,
+        mip_level_count: 1,
+        sample_count: 1,
+        view_formats: &[],
+    })
+}
+
+fn create_lighting_texture(device: &wgpu::Device, layers: u32, size: IVec2) -> wgpu::Texture {
+    device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("Lighting texture"),
+        size: wgpu::Extent3d {
+            width: size.x as u32,
+            height: size.y as u32,
+            depth_or_array_layers: layers,
+        },
+        usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
+        format: wgpu::TextureFormat::Rgba16Float,
+        dimension: wgpu::TextureDimension::D2,
+        mip_level_count: 1,
+        sample_count: 1,
+        view_formats: &[],
+    })
+}
+
+fn create_lighting_view(texture: &wgpu::Texture) -> wgpu::TextureView {
+    texture.create_view(&TextureViewDescriptor {
+        dimension: Some(wgpu::TextureViewDimension::D2Array),
+        array_layer_count: Some(LAYERS),
+        ..TextureViewDescriptor::default()
+    })
+}
+
+fn create_lighting_bind_group(
+    device: &wgpu::Device,
+    layout: &wgpu::BindGroupLayout,
+    scene_texture: &wgpu::Texture,
+    src: &wgpu::Texture,
+    dst: &wgpu::Texture,
+) -> wgpu::BindGroup {
+    device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: Some("Lighting bind group"),
+        layout,
+        entries: &[
+            wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&scene_texture.create_view(
+                    &TextureViewDescriptor {
+                        dimension: Some(wgpu::TextureViewDimension::D2),
+                        ..TextureViewDescriptor::default()
+                    },
+                )),
+            },
+            wgpu::BindGroupEntry {
+                binding: 1,
+                resource: wgpu::BindingResource::TextureView(&create_lighting_view(src)),
+            },
+            wgpu::BindGroupEntry {
+                binding: 2,
+                resource: wgpu::BindingResource::TextureView(&create_lighting_view(dst)),
+            },
+        ],
+    })
+}
+
+pub struct LightingRendererState {
+    lighting_pipeline: wgpu::ComputePipeline,
+    lighting_bind_group_layout: wgpu::BindGroupLayout,
+    uniform_buffer: wgpu::Buffer,
+    scene_texture: wgpu::Texture,
+    lighting_texture_a: wgpu::Texture,
+    lighting_texture_b: wgpu::Texture,
+    lighting_bind_group_a: wgpu::BindGroup,
+    lighting_bind_group_b: wgpu::BindGroup,
+}
+
+impl LightingRendererState {
+    pub fn lighting_view(&self) -> wgpu::TextureView {
+        create_lighting_view(&self.lighting_texture_a)
+    }
+
+    pub fn uniform_buffer(&self) -> &wgpu::Buffer {
+        &self.uniform_buffer
+    }
+
+    pub fn update_scene_texture(
+        &mut self,
+        device: &wgpu::Device,
+        queue: &wgpu::Queue,
+        sim: &mut SimManager,
+        camera: &Camera,
+    ) -> bool {
+        puffin::profile_function!();
+        let (lower, upper) = camera.viewport_bounds_world();
+        let origin = (lower * LIGHT_RESOLUTION).floor() / LIGHT_RESOLUTION;
+        let size = ((upper - origin) * LIGHT_RESOLUTION).ceil().as_ivec2();
+
+        let recreated = self.scene_texture.size().width != size.x as u32
+            || self.scene_texture.size().height != size.y as u32;
+
+        if recreated {
+            self.scene_texture = create_scene_texture(device, size);
+            self.lighting_texture_a = create_lighting_texture(device, LAYERS, size);
+            self.lighting_texture_b = create_lighting_texture(device, LAYERS, size);
+            self.lighting_bind_group_a = create_lighting_bind_group(
+                device,
+                &self.lighting_bind_group_layout,
+                &self.scene_texture,
+                &self.lighting_texture_a,
+                &self.lighting_texture_b,
+            );
+            self.lighting_bind_group_b = create_lighting_bind_group(
+                device,
+                &self.lighting_bind_group_layout,
+                &self.scene_texture,
+                &self.lighting_texture_b,
+                &self.lighting_texture_a,
+            );
+        }
+
+        queue.write_buffer(
+            &self.uniform_buffer,
+            0,
+            bytemuck::bytes_of(&LightingUniform {
+                origin: origin.to_array(),
+                _padding: [0.0; 2],
+            }),
+        );
+
+        let mut data: Vec<u8> = vec![
+            0;
+            (self.scene_texture.size().width * self.scene_texture.size().height * 4)
+                as usize
+        ];
+
+        for y in 0..self.scene_texture.size().height {
+            for x in 0..self.scene_texture.size().width {
+                let local = Vec2::new(x as f32, y as f32);
+                let world = origin + local * LIGHT_RESOLUTION;
+                let cell = sim
+                    .cell_manager
+                    .get_cell_from_game_position(world.x.round() as i32, world.y.round() as i32);
+                let material = cell.map_or(MaterialId::Void.def(), |c| c.material.def());
+
+                let idx = (x + y * self.scene_texture.size().width) as usize * 4;
+                data[idx] = material.emission.0;
+                data[idx + 1] = material.emission.1;
+                data[idx + 2] = material.emission.2;
+                data[idx + 3] = 0xFF - (material.opacity / 2);
+            }
+        }
+
+        queue.write_texture(
+            wgpu::TexelCopyTextureInfoBase {
+                texture: &self.scene_texture,
+                mip_level: 0,
+                origin: wgpu::Origin3d::ZERO,
+                aspect: wgpu::TextureAspect::All,
+            },
+            &data,
+            wgpu::TexelCopyBufferLayout {
+                bytes_per_row: Some(4 * self.scene_texture.size().width),
+                rows_per_image: Some(self.scene_texture.size().height),
+                offset: 0,
+            },
+            self.scene_texture.size(),
+        );
+
+        recreated
+    }
+
+    pub fn compute(&self, compute_pass: &mut wgpu::ComputePass) {
+        compute_pass.set_pipeline(&self.lighting_pipeline);
+        // must be even number of iterations so the result ends up in texture a
+        for i in 0..4 {
+            compute_pass.set_bind_group(
+                0,
+                if i % 2 == 0 {
+                    &self.lighting_bind_group_a
+                } else {
+                    &self.lighting_bind_group_b
+                },
+                &[],
+            );
+            compute_pass.dispatch_workgroups(
+                self.lighting_texture_a
+                    .size()
+                    .width
+                    .div_ceil(WORKGROUP_SIZE),
+                self.lighting_texture_a
+                    .size()
+                    .height
+                    .div_ceil(WORKGROUP_SIZE),
+                1,
+            );
+        }
+    }
+
+    pub fn new(device: &wgpu::Device) -> Self {
+        let lpv_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Light propagation volume shader"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("../shader/lpv.wgsl").into()),
+        });
+
+        let lighting_bind_group_layout =
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: None,
+                entries: &[
+                    // scene
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::Texture {
+                            multisampled: false,
+                            sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                            view_dimension: wgpu::TextureViewDimension::D2,
+                        },
+                        binding: 0,
+                        count: None,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                    },
+                    // src
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::Texture {
+                            multisampled: false,
+                            sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                            view_dimension: wgpu::TextureViewDimension::D2Array,
+                        },
+                        binding: 1,
+                        count: None,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                    },
+                    // dst
+                    wgpu::BindGroupLayoutEntry {
+                        ty: wgpu::BindingType::StorageTexture {
+                            access: wgpu::StorageTextureAccess::WriteOnly,
+                            format: wgpu::TextureFormat::Rgba16Float,
+                            view_dimension: wgpu::TextureViewDimension::D2Array,
+                        },
+                        binding: 2,
+                        count: None,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                    },
+                ],
+            });
+
+        let lighting_pipeline_layout =
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: None,
+                immediate_size: 0,
+                bind_group_layouts: &[Some(&lighting_bind_group_layout)],
+            });
+
+        let lighting_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&lighting_pipeline_layout),
+            entry_point: Some("propagate"),
+            compilation_options: wgpu::PipelineCompilationOptions::default(),
+            module: &lpv_shader,
+            cache: None,
+        });
+
+        let uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Lighting uniform buffer"),
+            size: size_of::<LightingUniform>() as u64,
+            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let size = IVec2::ONE;
+
+        let scene_texture = create_scene_texture(device, size);
+        let lighting_texture_a = create_lighting_texture(device, LAYERS, size);
+        let lighting_texture_b = create_lighting_texture(device, LAYERS, size);
+
+        let lighting_bind_group_a = create_lighting_bind_group(
+            device,
+            &lighting_bind_group_layout,
+            &scene_texture,
+            &lighting_texture_a,
+            &lighting_texture_b,
+        );
+        let lighting_bind_group_b = create_lighting_bind_group(
+            device,
+            &lighting_bind_group_layout,
+            &scene_texture,
+            &lighting_texture_b,
+            &lighting_texture_a,
+        );
+
+        LightingRendererState {
+            lighting_pipeline,
+            lighting_bind_group_layout,
+            uniform_buffer,
+            lighting_bind_group_a,
+            lighting_bind_group_b,
+            scene_texture,
+            lighting_texture_a,
+            lighting_texture_b,
+        }
+    }
+}
diff --git a/src/renderer/mod.rs b/src/renderer/mod.rs
index 3cf340d..8e94f55 100644
--- a/src/renderer/mod.rs
+++ b/src/renderer/mod.rs
@@ -1,5 +1,6 @@
 mod debug;
 mod instances;
+mod lighting;
 mod particles;
 mod ui;
 mod vfx;
@@ -15,7 +16,8 @@ use crate::{
     proc_gen::TilesetWorldGenerator,
     renderer::{
         debug::PhysicsDebugRendererState, instances::InstancesRendererState,
-        particles::ParticlesRendererState, ui::draw_egui, vfx::VfxRendererState,
+        lighting::LightingRendererState, particles::ParticlesRendererState, ui::draw_egui,
+        vfx::VfxRendererState,
     },
     sim::sim_manager::SimManager,
     vfx::VfxManager,
@@ -51,6 +53,7 @@ pub struct RendererState {
     particles_renderer_state: ParticlesRendererState,
     vfx_renderer_state: VfxRendererState,
     physics_debug_renderer_state: PhysicsDebugRendererState,
+    lighting_renderer_state: LightingRendererState,
 }
 
 impl RendererState {
@@ -150,8 +153,16 @@ impl RendererState {
             mapped_at_creation: false,
         });
 
-        let instances_renderer_state =
-            InstancesRendererState::new(&device, &config, &camera_uniform_buffer, &palette_buffer);
+        let lighting_renderer_state = LightingRendererState::new(&device);
+
+        let instances_renderer_state = InstancesRendererState::new(
+            &device,
+            &config,
+            &camera_uniform_buffer,
+            &palette_buffer,
+            &lighting_renderer_state.lighting_view(),
+            lighting_renderer_state.uniform_buffer(),
+        );
 
         let particles_renderer_state =
             ParticlesRendererState::new(&device, &config, &camera_uniform_buffer, &palette_buffer);
@@ -180,6 +191,7 @@ impl RendererState {
             particles_renderer_state,
             vfx_renderer_state,
             physics_debug_renderer_state,
+            lighting_renderer_state,
         }
     }
 
@@ -295,6 +307,16 @@ impl RendererState {
             bytemuck::bytes_of(&camera.to_uniform()),
         );
 
+        let lighting_recreated = self
+            .lighting_renderer_state
+            .update_scene_texture(&self.device, &self.queue, sim, camera);
+        if lighting_recreated {
+            self.instances_renderer_state.rebuild_bind_group(
+                &self.device,
+                &self.lighting_renderer_state.lighting_view(),
+                self.lighting_renderer_state.uniform_buffer(),
+            );
+        }
         self.instances_renderer_state
             .update_buffers(&self.device, &self.queue, sim, camera);
         self.particles_renderer_state
@@ -311,6 +333,17 @@ impl RendererState {
         }
 
         {
+            puffin::profile_scope!("Lighting compute pass");
+            let mut compute_pass: wgpu::ComputePass =
+                encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                    label: Some("Lighting compute pass"),
+                    timestamp_writes: None,
+                });
+
+            self.lighting_renderer_state.compute(&mut compute_pass);
+        }
+
+        {
             puffin::profile_scope!("Main render pass");
             let mut render_pass: wgpu::RenderPass<'_> =
                 encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
diff --git a/src/shader/instance.wgsl b/src/shader/instance.wgsl
index a549801..4b921ab 100644
--- a/src/shader/instance.wgsl
+++ b/src/shader/instance.wgsl
@@ -1,8 +1,16 @@
+const AMBIENT_LIGHT = vec4f(0.4, 0.4, 0.4, 1.0);
+const LIGHT_RESOLUTION = 1.0;
+
 struct Camera {
     scale: vec2f,
     centre: vec2f,
 };
 
+struct Lighting {
+    origin: vec2f,
+    _padding: vec2f,
+};
+
 struct Instance {
     centre: vec2f,
     cos_sin: vec2f,
@@ -15,6 +23,9 @@ struct Instance {
 @group(0) @binding(1) var<storage, read> palette: array<vec4f>;
 @group(0) @binding(2) var<storage, read> instances: array<Instance>;
 @group(0) @binding(3) var<storage, read> cells: array<u32>;
+@group(0) @binding(4) var lighting: texture_2d_array<f32>;
+@group(0) @binding(5) var light_sampler: sampler;
+@group(0) @binding(6) var<uniform> light_params: Lighting;
 
 struct VertexOutput {
     @builtin(position) clip_position: vec4f,
@@ -67,5 +78,14 @@ fn fs_main(in: VertexOutput) -> @location(0) vec4f {
         discard;
     }
 
-    return palette[material];
+    var total = vec3f(0.0);
+    let light_dims = vec2f(textureDimensions(lighting));
+    let light_texel = (in.world - light_params.origin) * LIGHT_RESOLUTION;
+    let light_uv = light_texel / light_dims;
+    for (var d = 0u; d < 4u; d++) {
+        total += textureSample(lighting, light_sampler, light_uv, d).rgb;
+    }
+    let light = vec4f(total.r, total.g, total.b, 1.0);
+
+    return palette[material] * (light + AMBIENT_LIGHT);
 }
diff --git a/src/shader/lpv.wgsl b/src/shader/lpv.wgsl
new file mode 100644
index 0000000..1f84e40
--- /dev/null
+++ b/src/shader/lpv.wgsl
@@ -0,0 +1,31 @@
+const DIRS = array<vec2i, 4>(vec2i(1, 0), vec2i(0, 1), vec2i(-1, 0), vec2i(0, -1));
+const ATTEN: f32 = 0.999;
+const W_FWD: f32 = 0.70;
+const W_SIDE: f32 = 0.15;
+
+@group(0) @binding(0) var scene: texture_2d<f32>;
+@group(0) @binding(1) var src: texture_2d_array<f32>;
+@group(0) @binding(2) var dst: texture_storage_2d_array<rgba16float, write>;
+
+@compute @workgroup_size(8, 8, 1)
+fn propagate(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let dims = textureDimensions(dst);
+    if gid.x >= dims.x || gid.y >= dims.y {
+        return;
+    }
+
+    let p = vec2<i32>(gid.xy);
+    let s = textureLoad(scene, p, 0);
+    let t = s.a;
+
+    for (var d = 0u; d < 4u; d++) {
+        let q = p - DIRS[d];
+        var inc = vec3f(0.0);
+        inc += textureLoad(src, q, d, 0).rgb * W_FWD;
+        inc += textureLoad(src, q, (d + 1u) % 4u, 0).rgb * W_SIDE;
+        inc += textureLoad(src, q, (d + 3u) % 4u, 0).rgb * W_SIDE;
+
+        let o = inc * t * ATTEN + s.rgb * 0.25;
+        textureStore(dst, p, d, vec4f(o, 1.0));
+    }
+}
\ No newline at end of file