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(1) var<storage, read> types: array<GrassType>;
@group(0) @binding(3) var<storage, read> blade_tiles: array<GrassTileEntry>;
@group(0) @binding(5) var<storage, read_write> draw_commands: GrassDrawCommands;
@group(0) @binding(6) var<storage, read_write> high_blades: array<GrassBlade>;
@group(0) @binding(7) var<storage, read_write> low_blades: array<GrassBlade>;
@group(0) @binding(8) var height_texture: texture_2d_array<f32>;

@compute @workgroup_size(256)
fn generate_blades(
    @builtin(workgroup_id) workgroup: vec3<u32>,
    @builtin(local_invocation_index) lane: u32,
) {
    let tile_index = workgroup.x / GRASS_WORKGROUPS_PER_TILE;
    let chunk = workgroup.x % GRASS_WORKGROUPS_PER_TILE;
    if tile_index >= min(atomicLoad(&draw_commands.blade_tile_count), GRASS_MAX_BLADE_TILES) {
        return;
    }
    let entry = blade_tiles[tile_index].data;
    let tile = entry.xy;
    let lod = u32(entry.z);
    let blade_index = chunk * 256u + lane;

    let stride = i32(1u << lod);
    let lane_cell = vec2<i32>(i32(blade_index % 80u), i32(blade_index / 80u));
    var base_cell = tile * 80 + lane_cell * stride;
    for (var level = lod; level >= 1u; level--) {
        let group = base_cell >> vec2<u32>(level, level);
        let pick_hash = grass_hash_tile(group, 7000u + level);
        base_cell += vec2<i32>(i32(pick_hash & 1u), i32((pick_hash >> 1u) & 1u))
            * (1 << (level - 1u));
    }

    var hash = grass_hash_tile(base_cell, 1u);
    let jitter_x = grass_hash01(hash);
    hash = grass_scramble(hash);
    let jitter_z = grass_hash01(hash);
    hash = grass_scramble(hash);
    var world_xz = (vec2<f32>(base_cell)
        + vec2<f32>(0.07, 0.07)
        + vec2<f32>(jitter_x, jitter_z) * 0.86)
        * (GRASS_TILE_SIZE / 80.0);

    let camera_xz = uniforms.camera_position.xz;
    let high_radius = uniforms.radii.x;
    let far_radius = uniforms.radii.z;
    let planar_distance = distance(world_xz, camera_xz);
    let lod_ground = uniforms.radii.w
        + grass_sample_height(
            height_texture,
            world_xz,
            uniforms.height_map,
            uniforms.height_bounds,
            planar_distance,
        );
    let blade_distance = distance(
        vec3<f32>(world_xz.x, lod_ground, world_xz.y),
        uniforms.camera_position.xyz,
    );

    if lod >= 1u
        && (grass_lod_for_distance(blade_distance, high_radius) != lod
            || blade_distance >= far_radius)
    {
        return;
    }
    if lod == 0u && blade_distance >= high_radius {
        return;
    }

    let parent_level = lod + 1u;
    let parent_group = base_cell >> vec2<u32>(parent_level, parent_level);
    let parent_hash = grass_hash_tile(parent_group, 7000u + parent_level);
    let quadrant = (base_cell >> vec2<u32>(lod, lod)) & vec2<i32>(1, 1);
    let promoted = i32(parent_hash & 1u) == quadrant.x
        && i32((parent_hash >> 1u) & 1u) == quadrant.y;

    var death_fade = 1.0;
    if !promoted || lod == GRASS_MAX_BLADE_LOD {
        var ring_start = high_radius * 0.78;
        var ring_end = high_radius;
        if lod >= 1u {
            ring_start = high_radius * f32(1u << (lod - 1u));
            ring_end = min(high_radius * f32(1u << lod), far_radius);
        }
        if lod == GRASS_MAX_BLADE_LOD {
            ring_end = far_radius;
        }
        let death_hash = grass_hash01(grass_hash_tile(base_cell, 9100u));
        let death_distance = mix(ring_start, ring_end, death_hash);
        if blade_distance > death_distance {
            return;
        }
        let fade_width = max((ring_end - ring_start) * 0.2, 1.0);
        death_fade = 1.0 - smoothstep(death_distance - fade_width, death_distance, blade_distance);
    }
    if !grass_inside_domain(world_xz, uniforms.counts.z, uniforms.domain) {
        return;
    }

    if uniforms.height_map.w > 1.5 {
        let texel = uniforms.height_map.z;
        let height_east = uniforms.radii.w
            + grass_sample_height(
                height_texture,
                world_xz + vec2<f32>(texel, 0.0),
                uniforms.height_map,
                uniforms.height_bounds,
                planar_distance,
            );
        let height_north = uniforms.radii.w
            + grass_sample_height(
                height_texture,
                world_xz + vec2<f32>(0.0, texel),
                uniforms.height_map,
                uniforms.height_bounds,
                planar_distance,
            );
        let slope_normal = normalize(vec3<f32>(
            lod_ground - height_east,
            texel,
            lod_ground - height_north,
        ));
        let slope = 1.0 - slope_normal.y;
        let splat_noise = terrain_fbm(world_xz * 0.05, 2u, 977u);
        let rock_amount = smoothstep(
            uniforms.splat.x - 0.12,
            uniforms.splat.x + 0.12,
            slope + splat_noise * 0.08,
        );
        let snow_amount = smoothstep(
            uniforms.splat.y - 14.0,
            uniforms.splat.y + 14.0,
            lod_ground + splat_noise * 18.0,
        );
        if grass_hash01(hash) < max(rock_amount, snow_amount) {
            return;
        }
        hash = grass_scramble(hash);
    }

    if lod == 0u {
        let keep = 1.0 - 0.75 * smoothstep(high_radius * 0.78, high_radius, blade_distance);
        if grass_hash01(hash) > keep {
            return;
        }
        hash = grass_scramble(hash);
    }

    let type_count = uniforms.counts.x;
    let type_index = grass_select_type(world_xz, uniforms.params.x, type_count, grass_hash01(hash));
    hash = grass_scramble(hash);
    let blade_type = types[type_index];

    if grass_hash01(hash) > blade_type.shape.w {
        return;
    }
    hash = grass_scramble(hash);

    let clump = grass_voronoi(world_xz, uniforms.params.y, 23u);
    var height = mix(blade_type.shape.x, blade_type.shape.y, grass_hash01(hash));
    hash = grass_scramble(hash);
    height *= 1.0 + (clump.cell_value - 0.5) * 2.0 * blade_type.misc.z;

    height *= death_fade;
    if height < 0.05 {
        return;
    }

    world_xz = mix(world_xz, clump.point, blade_type.misc.w * 0.4);

    let facing_angle = grass_hash01(hash) * GRASS_TAU;
    hash = grass_scramble(hash);
    var facing = vec2<f32>(cos(facing_angle), sin(facing_angle));
    let away_from_clump = normalize(world_xz - clump.point + vec2<f32>(1e-4, 0.0));
    facing = normalize(mix(facing, away_from_clump, blade_type.misc2.x));

    var tilt = mix(blade_type.curve.x, blade_type.curve.y, grass_hash01(hash));
    hash = grass_scramble(hash);
    let bend = mix(blade_type.curve.z, blade_type.curve.w, grass_hash01(hash));
    hash = grass_scramble(hash);

    let wind_strength = grass_wind_sample(world_xz, uniforms.params.w, uniforms.wind);

    let pusher_count = uniforms.counts.y;
    for (var pusher_index = 0u; pusher_index < pusher_count; pusher_index++) {
        let pusher = uniforms.pushers[pusher_index];
        let delta = world_xz - pusher.xz;
        let push_distance = length(delta);
        if push_distance < pusher.w && push_distance > 1e-4 {
            let strength = uniforms.pusher_strength[pusher_index / 4u][pusher_index % 4u]
                * (1.0 - push_distance / pusher.w);
            let influence = clamp(strength, 0.0, 1.0);
            facing = normalize(mix(facing, delta / push_distance, influence));
            tilt = min(tilt + strength * 0.9, 1.45);
            height *= 1.0 - 0.25 * influence;
        }
    }

    let relative = world_xz - vec2<f32>(uniforms.camera_tile.xy) * GRASS_TILE_SIZE;
    let ground = uniforms.radii.w
        + grass_sample_height(
            height_texture,
            world_xz,
            uniforms.height_map,
            uniforms.height_bounds,
            planar_distance,
        );
    var blade: GrassBlade;
    let width = blade_type.shape.z
        * clamp(sqrt(blade_distance / (high_radius * 0.78)), 1.0, 5.0);
    blade.position_height = vec4<f32>(relative.x, ground, relative.y, height);
    blade.facing_type = vec4<f32>(facing.x, facing.y, grass_hash01(hash), f32(type_index));
    blade.shape = vec4<f32>(tilt, bend, width, wind_strength);
    blade.extra = vec4<f32>(clump.cell_value, 0.0, 0.0, 0.0);

    if lod == 0u {
        let slot = atomicAdd(&draw_commands.high_instance_count, 1u);
        if slot >= GRASS_MAX_BLADES {
            atomicSub(&draw_commands.high_instance_count, 1u);
            return;
        }
        high_blades[slot] = blade;
    } else {
        let slot = atomicAdd(&draw_commands.low_instance_count, 1u);
        if slot >= GRASS_MAX_BLADES {
            atomicSub(&draw_commands.low_instance_count, 1u);
            return;
        }
        low_blades[slot] = blade;
    }
}
