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