lightingHEADmaster

3 files changed, 488 insertions, 24 deletions

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 {