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