struct GrassDispatch {
    generate_x: atomic<u32>,
    generate_y: u32,
    generate_z: u32,
    padding: u32,
};

struct GrassDrawCommands {
    high_index_count: u32,
    high_instance_count: atomic<u32>,
    high_first_index: u32,
    high_base_vertex: i32,
    high_first_instance: u32,
    low_index_count: u32,
    low_instance_count: atomic<u32>,
    low_first_index: u32,
    low_base_vertex: i32,
    low_first_instance: u32,
    far_vertex_count: u32,
    far_instance_count: atomic<u32>,
    far_first_vertex: u32,
    far_first_instance: u32,
    blade_tile_count: atomic<u32>,
    padding: u32,
};

@group(0) @binding(0) var<uniform> uniforms: GrassUniforms;
@group(0) @binding(2) var<storage, read_write> dispatch_args: GrassDispatch;
@group(0) @binding(3) var<storage, read_write> blade_tiles: array<GrassTileEntry>;
@group(0) @binding(4) var<storage, read_write> far_tiles: array<GrassTileEntry>;
@group(0) @binding(5) var<storage, read_write> draw_commands: GrassDrawCommands;

fn aabb_visible(min_corner: vec3<f32>, max_corner: 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(min_corner.x, max_corner.x, plane.x >= 0.0),
            select(min_corner.y, max_corner.y, plane.y >= 0.0),
            select(min_corner.z, max_corner.z, plane.z >= 0.0),
        );
        if dot(plane.xyz, positive) + plane.w < 0.0 {
            return false;
        }
    }
    return true;
}

fn tile_visible(origin: vec2<f32>, span: f32) -> bool {
    let ground = uniforms.radii.w;
    let height_low = ground + uniforms.height_bounds.x - 0.5;
    let height_high = ground + uniforms.height_bounds.y + 1.8;
    let min_corner = vec3<f32>(origin.x - 1.0, height_low, origin.y - 1.0);
    let max_corner = vec3<f32>(origin.x + span + 1.0, height_high, origin.y + span + 1.0);
    return aabb_visible(min_corner, max_corner);
}

fn tile_distance_range(origin: vec2<f32>, span: f32) -> vec2<f32> {
    let ground = uniforms.radii.w;
    let camera = uniforms.camera_position.xyz;
    let min_corner = vec3<f32>(origin.x, ground + uniforms.height_bounds.x - 0.5, origin.y);
    let max_corner = vec3<f32>(
        origin.x + span,
        ground + uniforms.height_bounds.y + 1.8,
        origin.y + span,
    );
    let closest = clamp(camera, min_corner, max_corner);
    let center = (min_corner + max_corner) * 0.5;
    let far_corner = vec3<f32>(
        select(min_corner.x, max_corner.x, camera.x < center.x),
        select(min_corner.y, max_corner.y, camera.y < center.y),
        select(min_corner.z, max_corner.z, camera.z < center.z),
    );
    return vec2<f32>(distance(camera, closest), distance(camera, far_corner));
}

fn emit_blade_tile(tile: vec2<i32>, lod: u32) {
    let slot = atomicAdd(&draw_commands.blade_tile_count, 1u);
    if slot >= GRASS_MAX_BLADE_TILES {
        atomicSub(&draw_commands.blade_tile_count, 1u);
        return;
    }
    blade_tiles[slot] = GrassTileEntry(vec4<i32>(tile.x, tile.y, i32(lod), 0));
    atomicAdd(&dispatch_args.generate_x, GRASS_WORKGROUPS_PER_TILE);
}

@compute @workgroup_size(64)
fn cull_tiles(@builtin(global_invocation_id) id: vec3<u32>) {
    let tiles_per_side = uniforms.camera_tile.z;
    let total = u32(tiles_per_side * tiles_per_side);
    if id.x >= total {
        return;
    }
    let half_span = uniforms.camera_tile.w;
    let local = vec2<i32>(i32(id.x) % tiles_per_side, i32(id.x) / tiles_per_side) - vec2<i32>(half_span, half_span);
    let tile = uniforms.camera_tile.xy + local;
    let camera_xz = uniforms.camera_position.xz;
    let high_radius = uniforms.radii.x;
    let low_radius = uniforms.radii.y;
    let far_radius = uniforms.radii.z;
    let finite = uniforms.counts.z;
    let origin = vec2<f32>(tile) * GRASS_TILE_SIZE;

    let tile_range = tile_distance_range(origin, GRASS_TILE_SIZE);
    if tile_range.x < high_radius
        && grass_rect_overlaps(origin, origin + vec2<f32>(GRASS_TILE_SIZE), finite, uniforms.domain)
        && tile_visible(origin, GRASS_TILE_SIZE)
    {
        emit_blade_tile(tile, 0u);
    }

    for (var lod = 1u; lod <= GRASS_MAX_BLADE_LOD; lod++) {
        let mask = i32((1u << lod) - 1u);
        if (tile.x & mask) != 0 || (tile.y & mask) != 0 {
            continue;
        }
        let span = GRASS_TILE_SIZE * f32(1u << lod);
        let group_range = tile_distance_range(origin, span);
        let ring_near = high_radius * f32(1u << (lod - 1u));
        var ring_far = high_radius * f32(1u << lod);
        if lod == GRASS_MAX_BLADE_LOD {
            ring_far = far_radius;
        }
        if group_range.y < ring_near || group_range.x >= min(ring_far, far_radius) {
            continue;
        }
        if grass_rect_overlaps(origin, origin + vec2<f32>(span), finite, uniforms.domain)
            && tile_visible(origin, span)
        {
            emit_blade_tile(tile, lod);
        }
    }

    if uniforms.counts.w == 1u && (tile.x & 3) == 0 && (tile.y & 3) == 0 {
        let center = origin + vec2<f32>(GRASS_FAR_TILE_SPAN * 0.5);
        let center_distance = distance(center, camera_xz);
        if center_distance < far_radius + GRASS_FAR_TILE_SPAN
            && grass_rect_overlaps(origin, origin + vec2<f32>(GRASS_FAR_TILE_SPAN), finite, uniforms.domain)
            && tile_visible(origin, GRASS_FAR_TILE_SPAN)
        {
            let slot = atomicAdd(&draw_commands.far_instance_count, 1u);
            if slot >= GRASS_MAX_FAR_TILES {
                atomicSub(&draw_commands.far_instance_count, 1u);
                return;
            }
            far_tiles[slot] = GrassTileEntry(vec4<i32>(tile.x, tile.y, 2, 0));
        }
    }
}
