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,
flip_winding: u32,
culling_mask: u32,
_pad_culling_0: u32,
_pad_culling_1: u32,
_pad_culling_2: u32,
};
struct MeshBounds {
center: vec3<f32>,
radius: 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>,
object_count: u32,
min_screen_pixel_size: f32,
projection_scale_y: f32,
camera_culling_mask: u32,
};
@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>;
struct MeshLodInfo {
lod_count: u32,
thresholds: array<f32, 3>,
};
@group(0) @binding(6)
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;
}
@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;
}
if (object.culling_mask & culling.camera_culling_mask) == 0u {
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 {
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;
}
}