use crate::ecs::camera::queries::query_active_camera_matrices;
use crate::ecs::mesh::components::Vertex;
use crate::ecs::prefab::resources::mesh_cache_iter;
use crate::ecs::world::{GLOBAL_TRANSFORM, SKIN, VISIBILITY, WATER, World};
use crate::render::wgpu::rendergraph::{PassExecutionContext, PassNode};
use std::collections::HashMap;

pub const INITIAL_WATER_VERTEX_BUFFER_SIZE: usize = 10_000;
pub const INITIAL_WATER_INDEX_BUFFER_SIZE: usize = 30_000;
pub const INITIAL_WATER_INSTANCES: usize = 256;
pub const INITIAL_WATER_MATERIALS: usize = 32;
pub const INITIAL_WATER_OBJECTS: usize = 256;
pub const INITIAL_WATER_BATCHES: usize = 32;
pub const BUFFER_GROWTH_FACTOR: f32 = 2.0;

const VOLUMETRIC_CUBE_MESH_NAME: &str = "__volumetric_cube__";

fn create_unit_cube_mesh() -> (Vec<Vertex>, Vec<u32>) {
    let half = 0.5;
    let vertices = vec![
        Vertex {
            position: [-half, -half, half],
            normal: [0.0, 0.0, 1.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, half],
            normal: [0.0, 0.0, 1.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, half],
            normal: [0.0, 0.0, 1.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, half],
            normal: [0.0, 0.0, 1.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, -half, -half],
            normal: [0.0, 0.0, -1.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [-1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, -half],
            normal: [0.0, 0.0, -1.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [-1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, -half],
            normal: [0.0, 0.0, -1.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [-1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, -half],
            normal: [0.0, 0.0, -1.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [-1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, -half],
            normal: [0.0, 1.0, 0.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, -half],
            normal: [0.0, 1.0, 0.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, half],
            normal: [0.0, 1.0, 0.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, half],
            normal: [0.0, 1.0, 0.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, -half, -half],
            normal: [0.0, -1.0, 0.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, -half],
            normal: [0.0, -1.0, 0.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, half],
            normal: [0.0, -1.0, 0.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, -half, half],
            normal: [0.0, -1.0, 0.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [1.0, 0.0, 0.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, -half],
            normal: [1.0, 0.0, 0.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, 1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, -half],
            normal: [1.0, 0.0, 0.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, 1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, half, half],
            normal: [1.0, 0.0, 0.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, 1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [half, -half, half],
            normal: [1.0, 0.0, 0.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, 1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, -half, -half],
            normal: [-1.0, 0.0, 0.0],
            tex_coords: [1.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, -1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, -half],
            normal: [-1.0, 0.0, 0.0],
            tex_coords: [1.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, -1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, half, half],
            normal: [-1.0, 0.0, 0.0],
            tex_coords: [0.0, 1.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, -1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
        Vertex {
            position: [-half, -half, half],
            normal: [-1.0, 0.0, 0.0],
            tex_coords: [0.0, 0.0],
            tex_coords_1: [0.0, 0.0],
            tangent: [0.0, 0.0, -1.0, 1.0],
            color: [1.0, 1.0, 1.0, 1.0],
        },
    ];
    let indices = vec![
        0, 1, 2, 0, 2, 3, 4, 6, 5, 4, 7, 6, 8, 9, 10, 8, 10, 11, 12, 14, 13, 12, 15, 14, 16, 17,
        18, 16, 18, 19, 20, 22, 21, 20, 23, 22,
    ];
    (vertices, indices)
}

fn compute_volumetric_bounds(volume_size: [f32; 3]) -> MeshBoundsData {
    let half_size = [
        volume_size[0] * 0.5,
        volume_size[1] * 0.5,
        volume_size[2] * 0.5,
    ];
    let radius =
        (half_size[0] * half_size[0] + half_size[1] * half_size[1] + half_size[2] * half_size[2])
            .sqrt();
    MeshBoundsData {
        center: [0.0, 0.0, 0.0],
        radius,
    }
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct WaterUniforms {
    view: [[f32; 4]; 4],
    projection: [[f32; 4]; 4],
    view_projection: [[f32; 4]; 4],
    camera_position: [f32; 4],
    time: f32,
    _padding: [f32; 3],
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable, PartialEq)]
struct WaterMaterialData {
    base_color: [f32; 4],
    water_color: [f32; 4],
    wave_height: f32,
    choppy: f32,
    speed: f32,
    freq: f32,
    specular_strength: f32,
    fresnel_power: f32,
    volume_shape: u32,
    volume_flow_type: u32,
    volume_size: [f32; 3],
    is_volumetric: u32,
    flow_direction: [f32; 2],
    flow_strength: f32,
    _flow_padding: f32,
}

impl std::hash::Hash for WaterMaterialData {
    fn hash<H: std::hash::Hasher>(&self, state: &mut H) {
        self.base_color.iter().for_each(|v| v.to_bits().hash(state));
        self.water_color
            .iter()
            .for_each(|v| v.to_bits().hash(state));
        self.wave_height.to_bits().hash(state);
        self.choppy.to_bits().hash(state);
        self.speed.to_bits().hash(state);
        self.freq.to_bits().hash(state);
        self.specular_strength.to_bits().hash(state);
        self.fresnel_power.to_bits().hash(state);
        self.volume_shape.hash(state);
        self.volume_flow_type.hash(state);
        self.volume_size
            .iter()
            .for_each(|v| v.to_bits().hash(state));
        self.is_volumetric.hash(state);
        self.flow_direction
            .iter()
            .for_each(|v| v.to_bits().hash(state));
        self.flow_strength.to_bits().hash(state);
    }
}

impl Eq for WaterMaterialData {}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct WaterInstance {
    model: [[f32; 4]; 4],
    material_index: u32,
    _padding: [u32; 3],
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct WaterObjectData {
    instance_index: u32,
    mesh_id: u32,
    batch_id: u32,
    _padding: u32,
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct MeshBoundsData {
    center: [f32; 3],
    radius: f32,
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct CullingUniforms {
    frustum_planes: [[f32; 4]; 6],
    object_count: u32,
    _padding: [u32; 3],
}

#[repr(C)]
#[derive(Copy, Clone, Debug, bytemuck::Pod, bytemuck::Zeroable)]
struct DrawIndexedIndirect {
    index_count: u32,
    instance_count: u32,
    first_index: u32,
    base_vertex: i32,
    first_instance: u32,
}

#[derive(Debug, Clone)]
struct WaterMeshData {
    vertex_offset: u32,
    index_offset: u32,
    index_count: u32,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
enum WaterType {
    Horizontal,
    Vertical,
    Volumetric,
}

pub struct WaterMeshPass {
    horizontal_pipeline: wgpu::RenderPipeline,
    vertical_pipeline: wgpu::RenderPipeline,
    volumetric_pipeline: wgpu::RenderPipeline,
    culling_pipeline: wgpu::ComputePipeline,

    uniform_buffer: wgpu::Buffer,
    culling_uniform_buffer: wgpu::Buffer,

    render_bind_group_layout: wgpu::BindGroupLayout,
    render_bind_group: wgpu::BindGroup,
    culling_bind_group_layout: wgpu::BindGroupLayout,
    culling_bind_group: Option<wgpu::BindGroup>,

    material_buffer: wgpu::Buffer,
    instance_buffer: wgpu::Buffer,
    object_buffer: wgpu::Buffer,
    mesh_bounds_buffer: wgpu::Buffer,
    visible_indices_buffer: wgpu::Buffer,
    indirect_buffer: wgpu::Buffer,
    indirect_reset_buffer: wgpu::Buffer,

    vertex_buffer: wgpu::Buffer,
    index_buffer: wgpu::Buffer,

    vertex_buffer_size: usize,
    index_buffer_size: usize,
    instance_buffer_size: usize,
    material_buffer_size: usize,
    object_buffer_size: usize,
    mesh_bounds_buffer_size: usize,
    visible_indices_buffer_size: usize,
    indirect_buffer_size: usize,

    current_vertex_offset: u32,
    current_index_offset: u32,

    meshes: HashMap<String, (u32, WaterMeshData)>,
    next_mesh_id: u32,

    horizontal_batches: Vec<usize>,
    vertical_batches: Vec<usize>,
    volumetric_batches: Vec<(usize, [f32; 3])>,
    batch_count: usize,
    object_count: u32,
    indirect_reset_count: usize,
    camera_position: [f32; 3],
}

fn extract_frustum_planes(view_projection: &nalgebra_glm::Mat4) -> [[f32; 4]; 6] {
    let m = view_projection;
    let row0 = [m[(0, 0)], m[(0, 1)], m[(0, 2)], m[(0, 3)]];
    let row1 = [m[(1, 0)], m[(1, 1)], m[(1, 2)], m[(1, 3)]];
    let row2 = [m[(2, 0)], m[(2, 1)], m[(2, 2)], m[(2, 3)]];
    let row3 = [m[(3, 0)], m[(3, 1)], m[(3, 2)], m[(3, 3)]];

    let planes = [
        [
            row3[0] + row0[0],
            row3[1] + row0[1],
            row3[2] + row0[2],
            row3[3] + row0[3],
        ],
        [
            row3[0] - row0[0],
            row3[1] - row0[1],
            row3[2] - row0[2],
            row3[3] - row0[3],
        ],
        [
            row3[0] + row1[0],
            row3[1] + row1[1],
            row3[2] + row1[2],
            row3[3] + row1[3],
        ],
        [
            row3[0] - row1[0],
            row3[1] - row1[1],
            row3[2] - row1[2],
            row3[3] - row1[3],
        ],
        [
            row3[0] + row2[0],
            row3[1] + row2[1],
            row3[2] + row2[2],
            row3[3] + row2[3],
        ],
        [
            row3[0] - row2[0],
            row3[1] - row2[1],
            row3[2] - row2[2],
            row3[3] - row2[3],
        ],
    ];

    planes.map(|p| {
        let len = (p[0] * p[0] + p[1] * p[1] + p[2] * p[2]).sqrt();
        if len > 0.0001 {
            [p[0] / len, p[1] / len, p[2] / len, p[3] / len]
        } else {
            p
        }
    })
}

impl WaterMeshPass {
    pub fn new(
        device: &wgpu::Device,
        color_format: wgpu::TextureFormat,
        depth_format: wgpu::TextureFormat,
    ) -> Self {
        let uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Uniform Buffer"),
            size: std::mem::size_of::<WaterUniforms>() as u64,
            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let culling_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Culling Uniform Buffer"),
            size: std::mem::size_of::<CullingUniforms>() as u64,
            usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let material_buffer_size = INITIAL_WATER_MATERIALS;
        let material_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Material Buffer"),
            size: (std::mem::size_of::<WaterMaterialData>() * material_buffer_size) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let instance_buffer_size = INITIAL_WATER_INSTANCES;
        let instance_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Instance Buffer"),
            size: (std::mem::size_of::<WaterInstance>() * instance_buffer_size) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let object_buffer_size = INITIAL_WATER_OBJECTS;
        let object_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Object Buffer"),
            size: (std::mem::size_of::<WaterObjectData>() * object_buffer_size) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let mesh_bounds_buffer_size = 64;
        let mesh_bounds_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Bounds Buffer"),
            size: (std::mem::size_of::<MeshBoundsData>() * mesh_bounds_buffer_size) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let visible_indices_buffer_size = INITIAL_WATER_INSTANCES;
        let visible_indices_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Visible Indices Buffer"),
            size: (std::mem::size_of::<u32>() * visible_indices_buffer_size) as u64,
            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let indirect_buffer_size = INITIAL_WATER_BATCHES;
        let indirect_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Indirect Buffer"),
            size: (std::mem::size_of::<DrawIndexedIndirect>() * indirect_buffer_size) as u64,
            usage: wgpu::BufferUsages::INDIRECT
                | wgpu::BufferUsages::STORAGE
                | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let indirect_reset_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Indirect Reset Buffer"),
            size: (std::mem::size_of::<DrawIndexedIndirect>() * indirect_buffer_size) as u64,
            usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let vertex_buffer_size = INITIAL_WATER_VERTEX_BUFFER_SIZE;
        let vertex_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Unified Vertex Buffer"),
            size: (std::mem::size_of::<Vertex>() * vertex_buffer_size) as u64,
            usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let index_buffer_size = INITIAL_WATER_INDEX_BUFFER_SIZE;
        let index_buffer = device.create_buffer(&wgpu::BufferDescriptor {
            label: Some("Water Mesh Unified Index Buffer"),
            size: (std::mem::size_of::<u32>() * index_buffer_size) as u64,
            usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::COPY_DST,
            mapped_at_creation: false,
        });

        let render_bind_group_layout =
            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
                label: Some("Water Mesh Render Bind Group Layout"),
                entries: &[
                    wgpu::BindGroupLayoutEntry {
                        binding: 0,
                        visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Uniform,
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 1,
                        visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 2,
                        visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 3,
                        visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                ],
            });

        let render_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("Water Mesh Render Bind Group"),
            layout: &render_bind_group_layout,
            entries: &[
                wgpu::BindGroupEntry {
                    binding: 0,
                    resource: uniform_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 1,
                    resource: material_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 2,
                    resource: instance_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 3,
                    resource: visible_indices_buffer.as_entire_binding(),
                },
            ],
        });

        let culling_bind_group_layout =
            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
                label: Some("Water Mesh Culling Bind Group Layout"),
                entries: &[
                    wgpu::BindGroupLayoutEntry {
                        binding: 0,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 1,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Uniform,
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 2,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 3,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 4,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: false },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 5,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: false },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 6,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                ],
            });

        let horizontal_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("Water Mesh Horizontal Shader"),
            source: wgpu::ShaderSource::Wgsl(include_str!("../../shaders/water_mesh.wgsl").into()),
        });

        let vertical_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("Water Mesh Vertical Shader"),
            source: wgpu::ShaderSource::Wgsl(
                include_str!("../../shaders/water_mesh_vertical.wgsl").into(),
            ),
        });

        let volumetric_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("Water Volume Shader"),
            source: wgpu::ShaderSource::Wgsl(
                include_str!("../../shaders/water_volume.wgsl").into(),
            ),
        });

        let culling_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("Water Mesh Culling Shader"),
            source: wgpu::ShaderSource::Wgsl(
                include_str!("../../shaders/water_mesh_culling.wgsl").into(),
            ),
        });

        let render_pipeline_layout =
            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("Water Mesh Render Pipeline Layout"),
                bind_group_layouts: &[&render_bind_group_layout],
                push_constant_ranges: &[],
            });

        let culling_pipeline_layout =
            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("Water Mesh Culling Pipeline Layout"),
                bind_group_layouts: &[&culling_bind_group_layout],
                push_constant_ranges: &[],
            });

        let vertex_layout = wgpu::VertexBufferLayout {
            array_stride: std::mem::size_of::<Vertex>() as u64,
            step_mode: wgpu::VertexStepMode::Vertex,
            attributes: &[
                wgpu::VertexAttribute {
                    format: wgpu::VertexFormat::Float32x3,
                    offset: 0,
                    shader_location: 0,
                },
                wgpu::VertexAttribute {
                    format: wgpu::VertexFormat::Float32x3,
                    offset: 12,
                    shader_location: 1,
                },
                wgpu::VertexAttribute {
                    format: wgpu::VertexFormat::Float32x2,
                    offset: 24,
                    shader_location: 2,
                },
            ],
        };

        let horizontal_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
            label: Some("Water Mesh Horizontal Pipeline"),
            layout: Some(&render_pipeline_layout),
            vertex: wgpu::VertexState {
                module: &horizontal_shader,
                entry_point: Some("vs_main"),
                buffers: std::slice::from_ref(&vertex_layout),
                compilation_options: Default::default(),
            },
            fragment: Some(wgpu::FragmentState {
                module: &horizontal_shader,
                entry_point: Some("fs_main"),
                targets: &[Some(wgpu::ColorTargetState {
                    format: color_format,
                    blend: Some(wgpu::BlendState::ALPHA_BLENDING),
                    write_mask: wgpu::ColorWrites::ALL,
                })],
                compilation_options: Default::default(),
            }),
            primitive: wgpu::PrimitiveState {
                topology: wgpu::PrimitiveTopology::TriangleList,
                front_face: wgpu::FrontFace::Ccw,
                cull_mode: None,
                ..Default::default()
            },
            depth_stencil: Some(wgpu::DepthStencilState {
                format: depth_format,
                depth_write_enabled: true,
                depth_compare: wgpu::CompareFunction::GreaterEqual,
                stencil: wgpu::StencilState::default(),
                bias: wgpu::DepthBiasState::default(),
            }),
            multisample: wgpu::MultisampleState::default(),
            multiview: None,
            cache: None,
        });

        let vertical_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
            label: Some("Water Mesh Vertical Pipeline"),
            layout: Some(&render_pipeline_layout),
            vertex: wgpu::VertexState {
                module: &vertical_shader,
                entry_point: Some("vs_main"),
                buffers: std::slice::from_ref(&vertex_layout),
                compilation_options: Default::default(),
            },
            fragment: Some(wgpu::FragmentState {
                module: &vertical_shader,
                entry_point: Some("fs_main"),
                targets: &[Some(wgpu::ColorTargetState {
                    format: color_format,
                    blend: Some(wgpu::BlendState::ALPHA_BLENDING),
                    write_mask: wgpu::ColorWrites::ALL,
                })],
                compilation_options: Default::default(),
            }),
            primitive: wgpu::PrimitiveState {
                topology: wgpu::PrimitiveTopology::TriangleList,
                front_face: wgpu::FrontFace::Ccw,
                cull_mode: None,
                ..Default::default()
            },
            depth_stencil: Some(wgpu::DepthStencilState {
                format: depth_format,
                depth_write_enabled: true,
                depth_compare: wgpu::CompareFunction::GreaterEqual,
                stencil: wgpu::StencilState::default(),
                bias: wgpu::DepthBiasState::default(),
            }),
            multisample: wgpu::MultisampleState::default(),
            multiview: None,
            cache: None,
        });

        let volumetric_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
            label: Some("Water Volume Pipeline"),
            layout: Some(&render_pipeline_layout),
            vertex: wgpu::VertexState {
                module: &volumetric_shader,
                entry_point: Some("vs_main"),
                buffers: &[vertex_layout],
                compilation_options: Default::default(),
            },
            fragment: Some(wgpu::FragmentState {
                module: &volumetric_shader,
                entry_point: Some("fs_main"),
                targets: &[Some(wgpu::ColorTargetState {
                    format: color_format,
                    blend: Some(wgpu::BlendState::ALPHA_BLENDING),
                    write_mask: wgpu::ColorWrites::ALL,
                })],
                compilation_options: Default::default(),
            }),
            primitive: wgpu::PrimitiveState {
                topology: wgpu::PrimitiveTopology::TriangleList,
                front_face: wgpu::FrontFace::Ccw,
                cull_mode: None,
                ..Default::default()
            },
            depth_stencil: Some(wgpu::DepthStencilState {
                format: depth_format,
                depth_write_enabled: false,
                depth_compare: wgpu::CompareFunction::GreaterEqual,
                stencil: wgpu::StencilState::default(),
                bias: wgpu::DepthBiasState::default(),
            }),
            multisample: wgpu::MultisampleState::default(),
            multiview: None,
            cache: None,
        });

        let culling_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
            label: Some("Water Mesh Culling Pipeline"),
            layout: Some(&culling_pipeline_layout),
            module: &culling_shader,
            entry_point: Some("main"),
            compilation_options: Default::default(),
            cache: None,
        });

        Self {
            horizontal_pipeline,
            vertical_pipeline,
            volumetric_pipeline,
            culling_pipeline,
            uniform_buffer,
            culling_uniform_buffer,
            render_bind_group_layout,
            render_bind_group,
            culling_bind_group_layout,
            culling_bind_group: None,
            material_buffer,
            instance_buffer,
            object_buffer,
            mesh_bounds_buffer,
            visible_indices_buffer,
            indirect_buffer,
            indirect_reset_buffer,
            vertex_buffer,
            index_buffer,
            vertex_buffer_size,
            index_buffer_size,
            instance_buffer_size,
            material_buffer_size,
            object_buffer_size,
            mesh_bounds_buffer_size,
            visible_indices_buffer_size,
            indirect_buffer_size,
            current_vertex_offset: 0,
            current_index_offset: 0,
            meshes: HashMap::new(),
            next_mesh_id: 0,
            horizontal_batches: Vec::new(),
            vertical_batches: Vec::new(),
            volumetric_batches: Vec::new(),
            batch_count: 0,
            object_count: 0,
            indirect_reset_count: 0,
            camera_position: [0.0, 0.0, 0.0],
        }
    }

    fn ensure_mesh_in_buffer(
        &mut self,
        device: &wgpu::Device,
        queue: &wgpu::Queue,
        name: &str,
        vertices: &[Vertex],
        indices: &[u32],
    ) -> Option<u32> {
        if let Some((mesh_id, _)) = self.meshes.get(name) {
            return Some(*mesh_id);
        }

        let vertex_count = vertices.len() as u32;
        let index_count = indices.len() as u32;

        let required_vertices = self.current_vertex_offset as usize + vertices.len();
        if required_vertices > self.vertex_buffer_size {
            let new_size = ((required_vertices as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.vertex_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Unified Vertex Buffer (Resized)"),
                size: (std::mem::size_of::<Vertex>() * new_size) as u64,
                usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.vertex_buffer_size = new_size;
            self.current_vertex_offset = 0;
            self.current_index_offset = 0;
            self.meshes.clear();
            self.next_mesh_id = 0;
            return None;
        }

        let required_indices = self.current_index_offset as usize + indices.len();
        if required_indices > self.index_buffer_size {
            let new_size = ((required_indices as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.index_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Unified Index Buffer (Resized)"),
                size: (std::mem::size_of::<u32>() * new_size) as u64,
                usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.index_buffer_size = new_size;
            self.current_vertex_offset = 0;
            self.current_index_offset = 0;
            self.meshes.clear();
            self.next_mesh_id = 0;
            return None;
        }

        let vertex_offset = self.current_vertex_offset;
        let index_offset = self.current_index_offset;

        queue.write_buffer(
            &self.vertex_buffer,
            (vertex_offset as usize * std::mem::size_of::<Vertex>()) as u64,
            bytemuck::cast_slice(vertices),
        );

        queue.write_buffer(
            &self.index_buffer,
            (index_offset as usize * std::mem::size_of::<u32>()) as u64,
            bytemuck::cast_slice(indices),
        );

        let bounds = Self::compute_mesh_bounds(vertices);
        let mesh_id = self.next_mesh_id;
        self.next_mesh_id += 1;

        if mesh_id as usize >= self.mesh_bounds_buffer_size {
            let new_size = ((self.mesh_bounds_buffer_size as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.mesh_bounds_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Bounds Buffer (Resized)"),
                size: (std::mem::size_of::<MeshBoundsData>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.mesh_bounds_buffer_size = new_size;
            self.culling_bind_group = None;
        }

        queue.write_buffer(
            &self.mesh_bounds_buffer,
            (mesh_id as usize * std::mem::size_of::<MeshBoundsData>()) as u64,
            bytemuck::cast_slice(&[bounds]),
        );

        self.meshes.insert(
            name.to_string(),
            (
                mesh_id,
                WaterMeshData {
                    vertex_offset,
                    index_offset,
                    index_count,
                },
            ),
        );

        self.current_vertex_offset += vertex_count;
        self.current_index_offset += index_count;

        Some(mesh_id)
    }

    fn compute_mesh_bounds(vertices: &[Vertex]) -> MeshBoundsData {
        if vertices.is_empty() {
            return MeshBoundsData {
                center: [0.0, 0.0, 0.0],
                radius: 0.0,
            };
        }

        let mut min = [f32::MAX, f32::MAX, f32::MAX];
        let mut max = [f32::MIN, f32::MIN, f32::MIN];

        for vertex in vertices {
            for index in 0..3 {
                min[index] = min[index].min(vertex.position[index]);
                max[index] = max[index].max(vertex.position[index]);
            }
        }

        let center = [
            (min[0] + max[0]) * 0.5,
            (min[1] + max[1]) * 0.5,
            (min[2] + max[2]) * 0.5,
        ];

        let mut radius = 0.0f32;
        for vertex in vertices {
            let dx = vertex.position[0] - center[0];
            let dy = vertex.position[1] - center[1];
            let dz = vertex.position[2] - center[2];
            let dist = (dx * dx + dy * dy + dz * dz).sqrt();
            radius = radius.max(dist);
        }

        MeshBoundsData { center, radius }
    }

    fn ensure_volumetric_mesh_with_bounds(
        &mut self,
        device: &wgpu::Device,
        queue: &wgpu::Queue,
        volume_size: [f32; 3],
    ) -> Option<u32> {
        if !self.meshes.contains_key(VOLUMETRIC_CUBE_MESH_NAME) {
            let (vertices, indices) = create_unit_cube_mesh();
            self.ensure_mesh_in_buffer(
                device,
                queue,
                VOLUMETRIC_CUBE_MESH_NAME,
                &vertices,
                &indices,
            );
        }

        let (_, base_mesh_data) = self.meshes.get(VOLUMETRIC_CUBE_MESH_NAME)?;
        let base_mesh_data = base_mesh_data.clone();

        let bounds = compute_volumetric_bounds(volume_size);
        let mesh_id = self.next_mesh_id;
        self.next_mesh_id += 1;

        if mesh_id as usize >= self.mesh_bounds_buffer_size {
            let new_size = ((self.mesh_bounds_buffer_size as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.mesh_bounds_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Bounds Buffer (Resized)"),
                size: (std::mem::size_of::<MeshBoundsData>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.mesh_bounds_buffer_size = new_size;
            self.culling_bind_group = None;
        }

        queue.write_buffer(
            &self.mesh_bounds_buffer,
            (mesh_id as usize * std::mem::size_of::<MeshBoundsData>()) as u64,
            bytemuck::cast_slice(&[bounds]),
        );

        self.meshes.insert(
            format!("__volumetric_{}__", mesh_id),
            (mesh_id, base_mesh_data),
        );

        Some(mesh_id)
    }

    fn rebuild_render_bind_group(&mut self, device: &wgpu::Device) {
        self.render_bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("Water Mesh Render Bind Group"),
            layout: &self.render_bind_group_layout,
            entries: &[
                wgpu::BindGroupEntry {
                    binding: 0,
                    resource: self.uniform_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 1,
                    resource: self.material_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 2,
                    resource: self.instance_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 3,
                    resource: self.visible_indices_buffer.as_entire_binding(),
                },
            ],
        });
    }

    fn rebuild_culling_bind_group(&mut self, device: &wgpu::Device) {
        self.culling_bind_group = Some(device.create_bind_group(&wgpu::BindGroupDescriptor {
            label: Some("Water Mesh Culling Bind Group"),
            layout: &self.culling_bind_group_layout,
            entries: &[
                wgpu::BindGroupEntry {
                    binding: 0,
                    resource: self.instance_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 1,
                    resource: self.culling_uniform_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 2,
                    resource: self.mesh_bounds_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 3,
                    resource: self.object_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 4,
                    resource: self.indirect_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 5,
                    resource: self.visible_indices_buffer.as_entire_binding(),
                },
                wgpu::BindGroupEntry {
                    binding: 6,
                    resource: self.material_buffer.as_entire_binding(),
                },
            ],
        }));
    }
}

impl PassNode<World> for WaterMeshPass {
    fn name(&self) -> &str {
        "water_mesh_pass"
    }

    fn reads(&self) -> Vec<&str> {
        vec![]
    }

    fn writes(&self) -> Vec<&str> {
        vec![]
    }

    fn reads_writes(&self) -> Vec<&str> {
        vec!["color", "depth"]
    }

    fn prepare(&mut self, device: &wgpu::Device, queue: &wgpu::Queue, world: &World) {
        self.horizontal_batches.clear();
        self.vertical_batches.clear();
        self.volumetric_batches.clear();
        self.batch_count = 0;
        self.object_count = 0;

        let Some(camera_matrices) = query_active_camera_matrices(world) else {
            return;
        };

        let view = camera_matrices.view;
        let projection = camera_matrices.projection;
        let view_projection = projection * view;

        let mut material_to_index: HashMap<WaterMaterialData, u32> = HashMap::new();
        let mut materials_data: Vec<WaterMaterialData> = Vec::new();
        let mut instances: Vec<WaterInstance> = Vec::new();
        let mut objects: Vec<WaterObjectData> = Vec::new();

        struct MeshBatchInfo {
            mesh_name: String,
            water_type: WaterType,
            instances: Vec<(u32, WaterInstance)>,
            volume_size: Option<[f32; 3]>,
        }
        let mut mesh_batches: Vec<MeshBatchInfo> = Vec::new();
        let mut needed_meshes: HashMap<String, (Vec<Vertex>, Vec<u32>)> = HashMap::new();

        for entity in world.query_entities(WATER | GLOBAL_TRANSFORM | VISIBILITY) {
            if world.entity_has_components(entity, SKIN) {
                continue;
            }
            let Some(visibility) = world.get_visibility(entity) else {
                continue;
            };
            if !visibility.visible {
                continue;
            }

            let Some(water) = world.get_water(entity) else {
                continue;
            };
            if !water.enabled {
                continue;
            }

            let Some(global_transform) = world.get_global_transform(entity) else {
                continue;
            };

            let mesh_name: String;
            if water.is_volumetric {
                mesh_name = VOLUMETRIC_CUBE_MESH_NAME.to_string();
                if !needed_meshes.contains_key(VOLUMETRIC_CUBE_MESH_NAME) {
                    let (vertices, indices) = create_unit_cube_mesh();
                    needed_meshes
                        .insert(VOLUMETRIC_CUBE_MESH_NAME.to_string(), (vertices, indices));
                }
            } else {
                let Some(render_mesh) = world.get_render_mesh(entity) else {
                    continue;
                };
                mesh_name = render_mesh.name.clone();

                if !needed_meshes.contains_key(&mesh_name) {
                    for (name, mesh) in mesh_cache_iter(&world.resources.mesh_cache) {
                        if name == &mesh_name {
                            needed_meshes.insert(
                                name.clone(),
                                (mesh.vertices.clone(), mesh.indices.clone()),
                            );
                            break;
                        }
                    }
                }

                if !needed_meshes.contains_key(&mesh_name) {
                    continue;
                }
            }

            let material_data = WaterMaterialData {
                base_color: water.base_color,
                water_color: water.water_color,
                wave_height: water.wave_height,
                choppy: water.choppy,
                speed: water.speed,
                freq: water.frequency,
                specular_strength: water.specular_strength,
                fresnel_power: water.fresnel_power,
                volume_shape: water.volume_shape as u32,
                volume_flow_type: water.volume_flow_type as u32,
                volume_size: water.volume_size,
                is_volumetric: if water.is_volumetric { 1 } else { 0 },
                flow_direction: water.flow_direction,
                flow_strength: water.flow_strength,
                _flow_padding: 0.0,
            };

            let material_index = *material_to_index.entry(material_data).or_insert_with(|| {
                let index = materials_data.len() as u32;
                materials_data.push(material_data);
                index
            });

            let model: [[f32; 4]; 4] = global_transform.0.into();
            let instance = WaterInstance {
                model,
                material_index,
                _padding: [0; 3],
            };

            let instance_index = instances.len() as u32;
            instances.push(instance);

            let water_type = if water.is_volumetric {
                WaterType::Volumetric
            } else if water.is_vertical {
                WaterType::Vertical
            } else {
                WaterType::Horizontal
            };

            if water_type == WaterType::Volumetric {
                mesh_batches.push(MeshBatchInfo {
                    mesh_name: mesh_name.clone(),
                    water_type,
                    instances: vec![(instance_index, instance)],
                    volume_size: Some(water.volume_size),
                });
            } else {
                let existing_batch = mesh_batches
                    .iter_mut()
                    .find(|b| b.mesh_name == mesh_name && b.water_type == water_type);

                if let Some(batch) = existing_batch {
                    batch.instances.push((instance_index, instance));
                } else {
                    mesh_batches.push(MeshBatchInfo {
                        mesh_name: mesh_name.clone(),
                        water_type,
                        instances: vec![(instance_index, instance)],
                        volume_size: None,
                    });
                }
            }
        }

        if instances.is_empty() {
            return;
        }

        for (name, (vertices, indices)) in &needed_meshes {
            if name != VOLUMETRIC_CUBE_MESH_NAME {
                self.ensure_mesh_in_buffer(device, queue, name, vertices, indices);
            }
        }

        let mut volumetric_mesh_ids: HashMap<usize, u32> = HashMap::new();
        for (batch_idx, batch_info) in mesh_batches.iter().enumerate() {
            if batch_info.water_type == WaterType::Volumetric
                && let Some(volume_size) = batch_info.volume_size
                && let Some(mesh_id) =
                    self.ensure_volumetric_mesh_with_bounds(device, queue, volume_size)
            {
                volumetric_mesh_ids.insert(batch_idx, mesh_id);
            }
        }

        let mut bind_groups_need_rebuild = false;

        if instances.len() > self.instance_buffer_size {
            let new_size = ((instances.len() as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.instance_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Instance Buffer (Resized)"),
                size: (std::mem::size_of::<WaterInstance>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.instance_buffer_size = new_size;
            bind_groups_need_rebuild = true;
        }

        if materials_data.len() > self.material_buffer_size {
            let new_size = ((materials_data.len() as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.material_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Material Buffer (Resized)"),
                size: (std::mem::size_of::<WaterMaterialData>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.material_buffer_size = new_size;
            bind_groups_need_rebuild = true;
        }

        for (batch_idx, batch_info) in mesh_batches.iter().enumerate() {
            let mesh_id = if batch_info.water_type == WaterType::Volumetric {
                let Some(&mesh_id) = volumetric_mesh_ids.get(&batch_idx) else {
                    continue;
                };
                mesh_id
            } else {
                let Some((mesh_id, _)) = self.meshes.get(&batch_info.mesh_name) else {
                    continue;
                };
                *mesh_id
            };

            let batch_id = self.batch_count as u32;

            for (instance_index, _) in &batch_info.instances {
                objects.push(WaterObjectData {
                    instance_index: *instance_index,
                    mesh_id,
                    batch_id,
                    _padding: 0,
                });
            }

            match batch_info.water_type {
                WaterType::Horizontal => self.horizontal_batches.push(self.batch_count),
                WaterType::Vertical => self.vertical_batches.push(self.batch_count),
                WaterType::Volumetric => {
                    let position = if let Some((idx, _)) = batch_info.instances.first() {
                        let inst = &instances[*idx as usize];
                        [inst.model[3][0], inst.model[3][1], inst.model[3][2]]
                    } else {
                        [0.0, 0.0, 0.0]
                    };
                    self.volumetric_batches.push((self.batch_count, position));
                }
            }

            self.batch_count += 1;
        }

        self.object_count = objects.len() as u32;

        if objects.len() > self.object_buffer_size {
            let new_size = ((objects.len() as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.object_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Object Buffer (Resized)"),
                size: (std::mem::size_of::<WaterObjectData>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.object_buffer_size = new_size;
            bind_groups_need_rebuild = true;
        }

        if instances.len() > self.visible_indices_buffer_size {
            let new_size = ((instances.len() as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.visible_indices_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Visible Indices Buffer (Resized)"),
                size: (std::mem::size_of::<u32>() * new_size) as u64,
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.visible_indices_buffer_size = new_size;
            bind_groups_need_rebuild = true;
        }

        if self.batch_count > self.indirect_buffer_size {
            let new_size = ((self.batch_count as f32) * BUFFER_GROWTH_FACTOR) as usize;
            self.indirect_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Indirect Buffer (Resized)"),
                size: (std::mem::size_of::<DrawIndexedIndirect>() * new_size) as u64,
                usage: wgpu::BufferUsages::INDIRECT
                    | wgpu::BufferUsages::STORAGE
                    | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.indirect_reset_buffer = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("Water Mesh Indirect Reset Buffer (Resized)"),
                size: (std::mem::size_of::<DrawIndexedIndirect>() * new_size) as u64,
                usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            self.indirect_buffer_size = new_size;
            bind_groups_need_rebuild = true;
        }

        if bind_groups_need_rebuild {
            self.rebuild_render_bind_group(device);
            self.culling_bind_group = None;
        }

        let view_inverse = nalgebra_glm::inverse(&view);
        let camera_position = [
            view_inverse[(0, 3)],
            view_inverse[(1, 3)],
            view_inverse[(2, 3)],
            1.0,
        ];
        self.camera_position = [camera_position[0], camera_position[1], camera_position[2]];

        self.volumetric_batches.sort_by(|a, b| {
            let dist_a = (a.1[0] - self.camera_position[0]).powi(2)
                + (a.1[1] - self.camera_position[1]).powi(2)
                + (a.1[2] - self.camera_position[2]).powi(2);
            let dist_b = (b.1[0] - self.camera_position[0]).powi(2)
                + (b.1[1] - self.camera_position[1]).powi(2)
                + (b.1[2] - self.camera_position[2]).powi(2);
            dist_b
                .partial_cmp(&dist_a)
                .unwrap_or(std::cmp::Ordering::Equal)
        });

        let time = world.resources.window.timing.uptime_milliseconds as f32 / 1000.0;

        let uniforms = WaterUniforms {
            view: view.into(),
            projection: projection.into(),
            view_projection: view_projection.into(),
            camera_position,
            time,
            _padding: [0.0; 3],
        };
        queue.write_buffer(&self.uniform_buffer, 0, bytemuck::cast_slice(&[uniforms]));

        let frustum_planes = extract_frustum_planes(&view_projection);
        let culling_uniforms = CullingUniforms {
            frustum_planes,
            object_count: self.object_count,
            _padding: [0; 3],
        };
        queue.write_buffer(
            &self.culling_uniform_buffer,
            0,
            bytemuck::cast_slice(&[culling_uniforms]),
        );

        queue.write_buffer(
            &self.material_buffer,
            0,
            bytemuck::cast_slice(&materials_data),
        );
        queue.write_buffer(&self.instance_buffer, 0, bytemuck::cast_slice(&instances));
        queue.write_buffer(&self.object_buffer, 0, bytemuck::cast_slice(&objects));

        let mut indirect_commands = Vec::with_capacity(self.batch_count);
        let mut indirect_reset_commands = Vec::with_capacity(self.batch_count);
        let mut visible_indices = Vec::with_capacity(instances.len());
        let mut first_instance_offset = 0u32;

        for (batch_idx, batch_info) in mesh_batches.iter().enumerate() {
            let mesh_data = if batch_info.water_type == WaterType::Volumetric {
                let Some(&mesh_id) = volumetric_mesh_ids.get(&batch_idx) else {
                    continue;
                };
                let mesh_name = format!("__volumetric_{}__", mesh_id);
                let Some((_, data)) = self.meshes.get(&mesh_name) else {
                    continue;
                };
                data.clone()
            } else {
                let Some((_, data)) = self.meshes.get(&batch_info.mesh_name) else {
                    continue;
                };
                data.clone()
            };

            let instance_count = batch_info.instances.len() as u32;

            indirect_commands.push(DrawIndexedIndirect {
                index_count: mesh_data.index_count,
                instance_count,
                first_index: mesh_data.index_offset,
                base_vertex: mesh_data.vertex_offset as i32,
                first_instance: first_instance_offset,
            });

            indirect_reset_commands.push(DrawIndexedIndirect {
                index_count: mesh_data.index_count,
                instance_count: 0,
                first_index: mesh_data.index_offset,
                base_vertex: mesh_data.vertex_offset as i32,
                first_instance: first_instance_offset,
            });

            for (instance_index, _) in &batch_info.instances {
                visible_indices.push(*instance_index);
            }

            first_instance_offset += instance_count;
        }

        queue.write_buffer(
            &self.indirect_buffer,
            0,
            bytemuck::cast_slice(&indirect_commands),
        );
        queue.write_buffer(
            &self.indirect_reset_buffer,
            0,
            bytemuck::cast_slice(&indirect_reset_commands),
        );
        queue.write_buffer(
            &self.visible_indices_buffer,
            0,
            bytemuck::cast_slice(&visible_indices),
        );

        self.indirect_reset_count = indirect_commands.len();

        if self.culling_bind_group.is_none() {
            self.rebuild_culling_bind_group(device);
        }
    }

    fn execute<'r, 'e>(
        &mut self,
        context: PassExecutionContext<'r, 'e, World>,
    ) -> crate::render::wgpu::rendergraph::Result<
        Vec<crate::render::wgpu::rendergraph::SubGraphRunCommand<'r>>,
    > {
        if self.batch_count == 0 || self.object_count == 0 {
            return Ok(context.into_sub_graph_commands());
        }

        let (color_view, color_load, color_store) = context.get_color_attachment("color")?;
        let (depth_view, depth_load, depth_store) = context.get_depth_attachment("depth")?;

        if let Some(ref culling_bind_group) = self.culling_bind_group {
            if self.indirect_reset_count > 0 {
                let copy_size =
                    (self.indirect_reset_count * std::mem::size_of::<DrawIndexedIndirect>()) as u64;
                context.encoder.copy_buffer_to_buffer(
                    &self.indirect_reset_buffer,
                    0,
                    &self.indirect_buffer,
                    0,
                    copy_size,
                );
            }

            {
                let mut compute_pass =
                    context
                        .encoder
                        .begin_compute_pass(&wgpu::ComputePassDescriptor {
                            label: Some("Water Mesh Culling Pass"),
                            timestamp_writes: None,
                        });

                compute_pass.set_pipeline(&self.culling_pipeline);
                compute_pass.set_bind_group(0, culling_bind_group, &[]);

                let workgroup_size = 64;
                let dispatch_size = self.object_count.div_ceil(workgroup_size);
                compute_pass.dispatch_workgroups(dispatch_size, 1, 1);
            }
        }

        {
            let mut render_pass = context
                .encoder
                .begin_render_pass(&wgpu::RenderPassDescriptor {
                    label: Some("Water Mesh Render Pass"),
                    color_attachments: &[Some(wgpu::RenderPassColorAttachment {
                        view: color_view,
                        resolve_target: None,
                        ops: wgpu::Operations {
                            load: color_load,
                            store: color_store,
                        },
                        depth_slice: None,
                    })],
                    depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
                        view: depth_view,
                        depth_ops: Some(wgpu::Operations {
                            load: depth_load,
                            store: depth_store,
                        }),
                        stencil_ops: None,
                    }),
                    timestamp_writes: None,
                    occlusion_query_set: None,
                });

            render_pass.set_bind_group(0, &self.render_bind_group, &[]);
            render_pass.set_vertex_buffer(0, self.vertex_buffer.slice(..));
            render_pass.set_index_buffer(self.index_buffer.slice(..), wgpu::IndexFormat::Uint32);

            if !self.horizontal_batches.is_empty() {
                render_pass.set_pipeline(&self.horizontal_pipeline);
                for &batch_index in &self.horizontal_batches {
                    let indirect_offset =
                        (batch_index * std::mem::size_of::<DrawIndexedIndirect>()) as u64;
                    render_pass.draw_indexed_indirect(&self.indirect_buffer, indirect_offset);
                }
            }

            if !self.vertical_batches.is_empty() {
                render_pass.set_pipeline(&self.vertical_pipeline);
                for &batch_index in &self.vertical_batches {
                    let indirect_offset =
                        (batch_index * std::mem::size_of::<DrawIndexedIndirect>()) as u64;
                    render_pass.draw_indexed_indirect(&self.indirect_buffer, indirect_offset);
                }
            }

            if !self.volumetric_batches.is_empty() {
                render_pass.set_pipeline(&self.volumetric_pipeline);
                for &(batch_index, _) in &self.volumetric_batches {
                    let indirect_offset =
                        (batch_index * std::mem::size_of::<DrawIndexedIndirect>()) as u64;
                    render_pass.draw_indexed_indirect(&self.indirect_buffer, indirect_offset);
                }
            }
        }

        Ok(context.into_sub_graph_commands())
    }
}
