struct ModelMatrix {
model: mat4x4<f32>,
normal_matrix: mat3x3<f32>,
};
struct ObjectData {
transform_index: u32,
mesh_id: u32,
material_id: u32,
batch_id: u32,
morph_weights: array<f32, 8>,
morph_target_count: u32,
morph_displacement_offset: u32,
mesh_vertex_offset: u32,
mesh_vertex_count: u32,
entity_id: u32,
is_overlay: u32,
skip_occlusion: u32,
_padding0: u32,
};
struct MeshBounds {
center: vec3<f32>,
radius: f32,
};
struct MeshBoundsAABB {
min: vec3<f32>,
_pad0: f32,
max: vec3<f32>,
_pad1: f32,
};
struct DrawIndexedIndirect {
index_count: u32,
instance_count: atomic<u32>,
first_index: u32,
base_vertex: i32,
first_instance: u32,
};
struct CullingUniforms {
frustum_planes: array<vec4<f32>, 6>,
view_projection: mat4x4<f32>,
screen_size: vec2<f32>,
hiz_mip_count: f32,
occlusion_enabled: u32,
object_count: u32,
min_screen_pixel_size: f32,
projection_scale_y: f32,
};
@group(0) @binding(0)
var<storage, read> transforms: array<ModelMatrix>;
@group(0) @binding(1)
var<storage, read> objects: array<ObjectData>;
@group(0) @binding(2)
var<uniform> culling: CullingUniforms;
@group(0) @binding(3)
var<storage, read> mesh_bounds: array<MeshBounds>;
@group(0) @binding(4)
var<storage, read_write> indirect_commands: array<DrawIndexedIndirect>;
@group(0) @binding(5)
var<storage, read_write> visible_indices: array<u32>;
@group(0) @binding(6)
var hiz_texture: texture_2d<f32>;
@group(0) @binding(7)
var<storage, read> mesh_aabbs: array<MeshBoundsAABB>;
struct MeshLodInfo {
lod_count: u32,
thresholds: array<f32, 3>,
};
@group(0) @binding(8)
var<storage, read> mesh_lod: array<MeshLodInfo>;
fn sphere_in_frustum(center: vec3<f32>, radius: f32) -> bool {
for (var i = 0u; i < 6u; i++) {
let plane = culling.frustum_planes[i];
let distance = dot(plane.xyz, center) + plane.w;
if distance < -radius {
return false;
}
}
return true;
}
struct ProjectedAABB {
screen_min: vec2<f32>,
screen_max: vec2<f32>,
nearest_z: f32,
valid: bool,
};
fn project_corners_to_screen(corners: array<vec3<f32>, 8>) -> ProjectedAABB {
var result: ProjectedAABB;
result.screen_min = vec2(1.0);
result.screen_max = vec2(0.0);
result.nearest_z = 0.0;
result.valid = true;
for (var i = 0u; i < 8u; i++) {
let clip = culling.view_projection * vec4(corners[i], 1.0);
if clip.w <= 0.0 {
result.valid = false;
return result;
}
let ndc = clip.xyz / clip.w;
let uv = vec2(ndc.x * 0.5 + 0.5, 1.0 - (ndc.y * 0.5 + 0.5));
result.screen_min = min(result.screen_min, uv);
result.screen_max = max(result.screen_max, uv);
result.nearest_z = max(result.nearest_z, ndc.z);
}
return result;
}
fn is_occluded(corners: array<vec3<f32>, 8>) -> bool {
if culling.occlusion_enabled == 0u {
return false;
}
let projected = project_corners_to_screen(corners);
if !projected.valid {
return false;
}
let pixel_pad = vec2(8.0) / culling.screen_size;
let padded_min = clamp(projected.screen_min - pixel_pad, vec2(0.0), vec2(1.0));
let padded_max = clamp(projected.screen_max + pixel_pad, vec2(0.0), vec2(1.0));
if padded_min.x >= padded_max.x || padded_min.y >= padded_max.y {
return false;
}
let rect_size = max(
(padded_max.x - padded_min.x) * culling.screen_size.x,
(padded_max.y - padded_min.y) * culling.screen_size.y
);
let mip = i32(clamp(ceil(log2(rect_size)), 0.0, culling.hiz_mip_count - 1.0));
let mip_size = vec2<f32>(textureDimensions(hiz_texture, mip));
let min_texel = clamp(vec2<i32>(floor(padded_min * mip_size)), vec2(0), vec2<i32>(mip_size) - vec2(1));
let max_texel = clamp(vec2<i32>(floor(padded_max * mip_size)), vec2(0), vec2<i32>(mip_size) - vec2(1));
var hiz_depth = 1.0;
for (var y = min_texel.y; y <= max_texel.y; y++) {
for (var x = min_texel.x; x <= max_texel.x; x++) {
let d = textureLoad(hiz_texture, vec2(x, y), mip).r;
hiz_depth = min(hiz_depth, d);
}
}
return projected.nearest_z < hiz_depth;
}
@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
let object_index = global_id.x;
if object_index >= culling.object_count {
return;
}
let object = objects[object_index];
if object.batch_id == 0xFFFFFFFFu {
return;
}
let bounds = mesh_bounds[object.mesh_id];
let transform = transforms[object.transform_index];
let local_center = vec4<f32>(bounds.center, 1.0);
let world_center = transform.model * local_center;
let world_pos = world_center.xyz;
let scale_x = length(vec3<f32>(transform.model[0][0], transform.model[0][1], transform.model[0][2]));
let scale_y = length(vec3<f32>(transform.model[1][0], transform.model[1][1], transform.model[1][2]));
let scale_z = length(vec3<f32>(transform.model[2][0], transform.model[2][1], transform.model[2][2]));
let max_scale = max(max(scale_x, scale_y), scale_z);
let bounding_radius = max_scale * bounds.radius;
var visible = sphere_in_frustum(world_pos, bounding_radius);
var screen_diameter = 0.0;
if visible {
let clip = culling.view_projection * vec4(world_pos, 1.0);
if clip.w > 0.0 {
screen_diameter = bounding_radius * culling.projection_scale_y * culling.screen_size.y / clip.w;
if culling.min_screen_pixel_size > 0.0 && screen_diameter < culling.min_screen_pixel_size {
visible = false;
}
}
}
if visible && culling.occlusion_enabled != 0u && object.is_overlay == 0u && object.skip_occlusion == 0u {
let aabb = mesh_aabbs[object.mesh_id];
let model = transforms[object.transform_index].model;
let edge0 = model[0].xyz * (aabb.max.x - aabb.min.x);
let edge1 = model[1].xyz * (aabb.max.y - aabb.min.y);
let edge2 = model[2].xyz * (aabb.max.z - aabb.min.z);
let base = (model * vec4(aabb.min, 1.0)).xyz;
let corners = array<vec3<f32>, 8>(
base,
base + edge0,
base + edge1,
base + edge0 + edge1,
base + edge2,
base + edge0 + edge2,
base + edge1 + edge2,
base + edge0 + edge1 + edge2,
);
if is_occluded(corners) {
visible = false;
}
}
if visible {
let lod_info = mesh_lod[object.mesh_id];
var lod_level = 0u;
if lod_info.lod_count > 1u {
lod_level = lod_info.lod_count - 1u;
for (var level = 0u; level < lod_info.lod_count - 1u; level++) {
if screen_diameter >= lod_info.thresholds[level] {
lod_level = level;
break;
}
}
}
let target_batch = object.batch_id + lod_level;
let write_index = atomicAdd(&indirect_commands[target_batch].instance_count, 1u);
let first_instance = indirect_commands[target_batch].first_instance;
visible_indices[first_instance + write_index] = object_index;
}
}