struct TerrainRenderUniforms {
    view_projection: mat4x4<f32>,
    camera_position: vec4<f32>,
    sun_direction: vec4<f32>,
    sun_color: vec4<f32>,
    ambient: vec4<f32>,
    params: vec4<f32>,
    splat: vec4<f32>,
    rock_color: vec4<f32>,
    snow_color: vec4<f32>,
    wind_time: vec4<f32>,
    frustum_planes: array<vec4<f32>, 6>,
    level_regions: array<vec4<f32>, 12>,
    cascade_view_projections: array<mat4x4<f32>, 4>,
    cascade_splits: vec4<f32>,
    cascade_atlas_offsets: array<vec4<f32>, 4>,
    cascade_atlas_scale: vec4<f32>,
};

struct TerrainInstance {
    data: vec4<f32>,
};

struct TerrainDrawCommand {
    index_count: u32,
    instance_count: atomic<u32>,
    first_index: u32,
    base_vertex: i32,
    first_instance: u32,
};

@group(0) @binding(0) var<uniform> uniforms: TerrainRenderUniforms;
@group(0) @binding(1) var<storage, read_write> visible_instances: array<TerrainInstance>;
@group(0) @binding(2) var<storage, read_write> draw_command: TerrainDrawCommand;

const TERRAIN_BLOCKS_PER_LEVEL: u32 = 36u;
const TERRAIN_MAX_VISIBLE: u32 = 512u;

fn block_visible(minimum: vec3<f32>, maximum: vec3<f32>) -> bool {
    for (var plane_index = 0u; plane_index < 6u; plane_index++) {
        let plane = uniforms.frustum_planes[plane_index];
        let positive = vec3<f32>(
            select(minimum.x, maximum.x, plane.x >= 0.0),
            select(minimum.y, maximum.y, plane.y >= 0.0),
            select(minimum.z, maximum.z, plane.z >= 0.0),
        );
        if dot(plane.xyz, positive) + plane.w < 0.0 {
            return false;
        }
    }
    return true;
}

@compute @workgroup_size(64)
fn cull_blocks(@builtin(global_invocation_id) id: vec3<u32>) {
    let level_count = u32(uniforms.params.y);
    if id.x >= level_count * TERRAIN_BLOCKS_PER_LEVEL {
        return;
    }
    let level = id.x / TERRAIN_BLOCKS_PER_LEVEL;
    let block_index = id.x % TERRAIN_BLOCKS_PER_LEVEL;
    let block = vec2<f32>(f32(block_index % 6u), f32(block_index / 6u));

    let block_world = uniforms.params.x * f32(1u << level) * 32.0;
    let region = uniforms.level_regions[level];
    let region_origin = region.xy - vec2<f32>(region.z, region.z);
    let origin = region_origin + block * block_world;

    if level > 0u {
        let finer = uniforms.level_regions[level - 1u];
        let finer_origin = finer.xy - vec2<f32>(finer.z, finer.z);
        let finer_end = finer.xy + vec2<f32>(finer.z, finer.z);
        let contained = origin.x >= finer_origin.x - 0.5
            && origin.y >= finer_origin.y - 0.5
            && origin.x + block_world <= finer_end.x + 0.5
            && origin.y + block_world <= finer_end.y + 0.5;
        if contained {
            return;
        }
    }

    let minimum = vec3<f32>(origin.x, uniforms.params.z, origin.y);
    let maximum = vec3<f32>(origin.x + block_world, uniforms.params.w, origin.y + block_world);
    if !block_visible(minimum, maximum) {
        return;
    }

    let slot = atomicAdd(&draw_command.instance_count, 1u);
    if slot >= TERRAIN_MAX_VISIBLE {
        atomicSub(&draw_command.instance_count, 1u);
        return;
    }
    visible_instances[slot] = TerrainInstance(vec4<f32>(origin.x, origin.y, 0.0, f32(level)));
}
