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));
}
}
}