Search across packages, versions, files, declarations, and comments. Use multiple terms to narrow results (AND), and add @vertex, @fragment, or @compute to filter by shader stage.
1//! The shader that performs the cluster-object intersection tests and assigns
2//! objects to clusters as appropriate.
3//!
4//! Assuming the froxel grid has size WxHxD, this shader is expected to run on a
5//! viewport of size WxH. It draws each *Z slice* as an axis aligned quad such
6//! that each fragment shader invocation represents a single cluster-object pair
7//! in the froxel grid. The fragment shader performs a finer test to see if the
8//! object might intersects the cluster and, if it succeeds, records the result.
9//! Because the result is written to storage buffers, color writes are disabled
10//! for this shader invocation.
11//!
12//! This shader runs twice: once to accumulate the *count* of each object type in
13//! each cluster, and once to *populate* the actual IDs. Both invocations of the
14//! shader must compute the exact same visibility results.
16import package::cluster::{
17 Aabb, CLUSTERABLE_OBJECT_TYPE_DECAL, CLUSTERABLE_OBJECT_TYPE_IRRADIANCE_VOLUME,
18 CLUSTERABLE_OBJECT_TYPE_POINT_LIGHT, CLUSTERABLE_OBJECT_TYPE_REFLECTION_PROBE,
19 CLUSTERABLE_OBJECT_TYPE_SPOT_LIGHT, CLUSTERABLE_OBJECT_TYPE_RECT_LIGHT, ClusterableObjectZSlice,
20 calculate_sphere_cluster_bounds, compute_view_from_world_scale
21};
22import package::render::clustered_forward;
23import package::light_probe::transpose_affine_matrix;
24import package::render::mesh_view_types::{
25 ClusterOffsetsAndCounts, ClusterableObjectIndexLists, ClusteredDecals, ClusteredLights,
26 LightProbes, Lights, POINT_LIGHT_FLAGS_SPOT_LIGHT_Y_NEGATIVE
27};
28import bevy_render::view::View;
29import bevy_render::maths::quat_rotate;
31struct Vertex {
32 @builtin(instance_index) instance_id: u32,
33 @location(0) position: vec2<f32>,
36/// Data output from the vertex shader and input to the fragment shader.
37struct Varyings {
38 @builtin(position) position: vec4<f32>,
39 /// The index of the Z slice we're rasterizing.
40 @location(0) @interpolate(flat) instance_id: u32,
41 /// The view-space center of the bounding sphere of the object.
42 @location(1) @interpolate(flat) sphere_position: vec3<f32>,
43 /// The view-space radius of the bounding sphere of the object.
44 @location(2) @interpolate(flat) sphere_radius: f32,
47/// The same as the `ClusterOffsetsAndCounts` structure, but with atomic fields
48/// so that we can write to it.
49struct ClusterOffsetsAndCountsAtomic {
50 data: array<ClusterOffsetsAndCountsElementAtomic>,
53/// The same as the `ClusterOffsetsAndCountsElement` structure, but with atomic
54/// fields so that we can write to it.
55struct ClusterOffsetsAndCountsElementAtomic {
56 offset: atomic<u32>,
57 point_lights: atomic<u32>,
58 spot_lights: atomic<u32>,
59 rect_lights: atomic<u32>,
60 reflection_probes: atomic<u32>,
61 irradiance_volumes: atomic<u32>,
62 decals: atomic<u32>,
63 pad_a: u32,
66/// The list of clusterable object Z slices that we read from to.
67@group(0) @binding(0) var<storage> z_slices: array<ClusterableObjectZSlice>;
68@if(!VERTEX_SHADER)
69/// The list of indices per cluster that we write to in the populate pass.
70@group(0) @binding(1) var<storage, read_write> index_lists: ClusterableObjectIndexLists;
71/// Information about each light.
72@group(0) @binding(2) var<storage> clustered_lights: ClusteredLights;
73/// Information about each light probe (reflection probe or irradiance volume).
74@group(0) @binding(3) var<uniform> light_probes: LightProbes;
75/// Information about each clustered decal.
76@group(0) @binding(4) var<storage> clustered_decals: ClusteredDecals;
77/// Information about the clusters as a whole, including the dimensions of the
78/// cluster grid.
79@group(0) @binding(5) var<uniform> lights: Lights;
80/// Information about the view.
81@group(0) @binding(6) var<uniform> view: View;
82@if(!VERTEX_SHADER && POPULATE_PASS) {
83/// The number of objects in each cluster, and the offset of each list.
84@group(0) @binding(7) var<storage> offsets_and_counts: ClusterOffsetsAndCounts;
85/// For each cluster, the counts of objects written *so far* to it.
86///
87/// We use this during the populate phase in order to write the ID of each object
88/// to the correct spot.
89@group(0) @binding(8) var<storage, read_write> scratchpad_offsets_and_counts:
90 ClusterOffsetsAndCountsAtomic;
92@if(!VERTEX_SHADER && !POPULATE_PASS)
93/// The number of objects in each cluster.
94///
95/// During the count pass, we write to this.
96@group(0) @binding(7) var<storage, read_write> offsets_and_counts: ClusterOffsetsAndCountsAtomic;
98/// The vertex entry point.
99@vertex
100fn vertex_main(vertex: Vertex) -> Varyings {
101 let instance_id = vertex.instance_id;
102 let object_index = z_slices[instance_id].object_index;
103 let object_type = z_slices[instance_id].object_type;
105 // Look up the world space bounding sphere of the object.
106 let bounding_sphere = get_object_bounding_sphere(object_index, object_type);
107 let position = bounding_sphere.xyz;
108 let radius = bounding_sphere.w;
110 let view_from_world_scale = compute_view_from_world_scale(view.world_from_view);
111 let max_view_from_world_scale = max(view_from_world_scale.x,
112 max(view_from_world_scale.y, view_from_world_scale.z));
113 let is_orthographic = view.clip_from_view[3].w == 1.0;
115 // Calculate an approximate AABB of the cluster by computing its bounding
116 // sphere and converting that to the AABB.
117 // It's possible to do better, as the CPU version of
118 // `assign_objects_to_clusters` does with its *iterative sphere refinement*
119 // algorithm. However, that's sequential. I believe that this simple
120 // approach is the standard way to do cluster assignment on GPU.
121 let cluster_bounds = calculate_sphere_cluster_bounds(
122 position,
123 radius,
124 view.view_from_world,
125 view.clip_from_view,
126 view_from_world_scale,
127 lights.cluster_dimensions.xyz,
128 lights.cluster_factors.zw,
129 is_orthographic
130 );
131 let cluster_bounds_xy = vec4<u32>(cluster_bounds.min.xy, cluster_bounds.max.xy + vec2<u32>(1u));
133 // Calculate the bounding sphere's center and radius in view space.
134 let view_position = (view.view_from_world * vec4(position, 1.0)).xyz;
135 let view_radius = max_view_from_world_scale * radius;
137 return Varyings(
138 calculate_vertex_position(vertex, cluster_bounds_xy),
139 instance_id,
140 view_position,
141 view_radius
142 );
145/// Returns the position of the quad vertex necessary to enclose all the
146/// fragments that represent the cluster AABB.
147/// The cluster bounds are supplied as `vec4(min X, min Y, max X, max Y)`.
148fn calculate_vertex_position(vertex: Vertex, cluster_bounds: vec4<u32>) -> vec4<f32> {
149 let framebuffer_position = vec2<u32>(
150 select(cluster_bounds.x, cluster_bounds.z, vertex.position.x == 1.0),
151 select(cluster_bounds.y, cluster_bounds.w, vertex.position.y == 1.0)
152 );
153 let vertex_position =
154 vec2<f32>(framebuffer_position) / vec2<f32>(lights.cluster_dimensions.xy);
155 return vec4(mix(vec2(-1.0, 1.0), vec2(1.0, -1.0), vertex_position), 0.0, 1.0);
158@if(!VERTEX_SHADER)
160/// Performs a fine-grained test to ensure that the object intersects a single
161/// froxel and records the result.
162@fragment
163fn fragment_main(varyings: Varyings) -> @location(0) vec4<f32> {
164 let instance_id = varyings.instance_id;
165 let object_index = z_slices[instance_id].object_index;
166 let object_type = z_slices[instance_id].object_type;
167 let z_slice = z_slices[instance_id].z_slice;
169 let is_orthographic = view.clip_from_view[3].w == 1.0;
170 let screen_size = view.viewport.zw;
171 let tile_size = screen_size / vec2<f32>(lights.cluster_dimensions.xy);
173 let z_near_far = compute_z_near_and_z_far(is_orthographic);
174 let z_near = z_near_far.x;
175 let z_far = z_near_far.y;
177 // Determine the AABB of the cluster this fragment represents.
178 let cluster_position = vec3<u32>(vec2<u32>(floor(varyings.position.xy)), z_slice);
179 let cluster_aabb = compute_aabb_for_cluster(
180 z_near,
181 z_far,
182 tile_size,
183 screen_size,
184 view.view_from_clip,
185 is_orthographic,
186 lights.cluster_dimensions.xyz,
187 cluster_position
188 );
189 let cluster_aabb_center = (cluster_aabb.max + cluster_aabb.min) * 0.5;
190 let cluster_aabb_half_size = (cluster_aabb.max - cluster_aabb.min) * 0.5;
192 // See if the object sphere intersects the AABB. If it doesn't, cull the
193 // object.
194 let object_intersects_cluster_aabb = sphere_intersects_aabb(
195 varyings.sphere_position,
196 varyings.sphere_radius,
197 cluster_aabb_center,
198 cluster_aabb_half_size
199 );
200 if (!object_intersects_cluster_aabb) {
201 return vec4<f32>(0.0);
202 }
204 // Do further, more precise culling for spot lights.
205 if (object_type == CLUSTERABLE_OBJECT_TYPE_SPOT_LIGHT && cull_spot_light(
206 object_index,
207 cluster_aabb_center,
208 length(cluster_aabb_half_size),
209 varyings.sphere_position,
210 varyings.sphere_radius
211 )) {
212 return vec4<f32>(0.0);
213 }
215 // Do further, more precise culling for rect lights.
216 if (object_type == CLUSTERABLE_OBJECT_TYPE_RECT_LIGHT && cull_rect_light(
217 object_index,
218 cluster_aabb_center,
219 length(cluster_aabb_half_size),
220 varyings.sphere_position
221 )) {
222 return vec4<f32>(0.0);
223 }
225 let cluster_index =
226 clustered_forward::fragment_cluster_index(cluster_position, lights.cluster_dimensions);
228 // If this is the populate pass, reserve a slot and write in the actual
229 // object index. Otherwise, if this is the count pass, just bump the
230 // appropriate counter.
231 @if(POPULATE_PASS)
232 let output_index = allocate_list_entry(cluster_index, object_type);
233 @if(POPULATE_PASS)
234 if (output_index < arrayLength(&index_lists.data)) {
235 index_lists.data[output_index] = object_index;
236 }
237 @else // POPULATE_PASS
238 increment_object_count(cluster_index, object_type);
240 return vec4<f32>(0.0);
243/// Returns true if the given sphere intersects the AABB with the given
244/// boundaries.
245fn sphere_intersects_aabb(
246 sphere_center: vec3<f32>,
247 sphere_radius: f32,
248 aabb_center: vec3<f32>,
249 aabb_half_size: vec3<f32>
250) -> bool {
251 let delta = max(vec3(0.0), abs(aabb_center - sphere_center) - aabb_half_size);
252 let dist_sq = dot(delta, delta);
253 return dist_sq <= sphere_radius * sphere_radius;
256/// See `bevy_light::cluster::assign::compute_aabb_for_cluster`.
257fn compute_aabb_for_cluster(
258 z_near: f32,
259 z_far: f32,
260 tile_size: vec2<f32>,
261 screen_size: vec2<f32>,
262 view_from_clip: mat4x4<f32>,
263 is_orthographic: bool,
264 cluster_dimensions: vec3<u32>,
265 ijk_u: vec3<u32>
266) -> Aabb {
267 let ijk = vec3<f32>(ijk_u);
269 // Calculate the minimum and maximum points in screen space
270 let p_min = ijk.xy * tile_size;
271 let p_max = p_min + tile_size;
273 var cluster_min: vec3<f32>;
274 var cluster_max: vec3<f32>;
275 if (is_orthographic) {
276 // Use linear depth slicing for orthographic
278 // Convert to view space at the cluster near and far planes
279 // NOTE: 1.0 is the near plane due to using reverse z projections
280 var p_min = screen_to_view(screen_size, view_from_clip, p_min, 0.0).xyz;
281 var p_max = screen_to_view(screen_size, view_from_clip, p_max, 0.0).xyz;
283 // calculate cluster depth using z_near and z_far
284 p_min.z = -z_near + (z_near - z_far) * ijk.z / f32(cluster_dimensions.z);
285 p_max.z = -z_near + (z_near - z_far) * (ijk.z + 1.0) / f32(cluster_dimensions.z);
287 cluster_min = min(p_min, p_max);
288 cluster_max = max(p_min, p_max);
289 } else {
290 // Convert to view space at the near plane
291 // NOTE: 1.0 is the near plane due to using reverse z projections
292 let p_min = screen_to_view(screen_size, view_from_clip, p_min, 1.0);
293 let p_max = screen_to_view(screen_size, view_from_clip, p_max, 1.0);
295 let z_far_over_z_near = -z_far / -z_near;
296 var cluster_near = 0.0;
297 if (ijk.z != 0.0) {
298 cluster_near = -z_near *
299 pow(z_far_over_z_near, (ijk.z - 1.0) / f32(cluster_dimensions.z - 1u));
300 }
301 // NOTE: This could be simplified to:
302 // cluster_far = cluster_near * z_far_over_z_near;
303 var cluster_far: f32;
304 if (cluster_dimensions.z == 1u) {
305 cluster_far = -z_far;
306 } else {
307 cluster_far = -z_near * pow(z_far_over_z_near, ijk.z / f32(cluster_dimensions.z - 1u));
308 }
310 // Calculate the four intersection points of the min and max points with the cluster near and far planes
311 let p_min_near = line_intersection_to_z_plane(vec3(0.0), p_min.xyz, cluster_near);
312 let p_min_far = line_intersection_to_z_plane(vec3(0.0), p_min.xyz, cluster_far);
313 let p_max_near = line_intersection_to_z_plane(vec3(0.0), p_max.xyz, cluster_near);
314 let p_max_far = line_intersection_to_z_plane(vec3(0.0), p_max.xyz, cluster_far);
316 cluster_min = min(min(p_min_near, p_min_far), min(p_max_near, p_max_far));
317 cluster_max = max(max(p_min_near, p_min_far), max(p_max_near, p_max_far));
318 }
320 return Aabb(cluster_min, cluster_max);
323/// Converts a screen-space position to a view-space position.
324/// See `bevy_light::cluster::assign::screen_to_view`.
325fn screen_to_view(
326 screen_size: vec2<f32>,
327 view_from_clip: mat4x4<f32>,
328 screen: vec2<f32>,
329 ndc_z: f32
330) -> vec4<f32> {
331 let tex_coord = screen / screen_size;
332 let clip = vec4(
333 tex_coord.x * 2.0 - 1.0,
334 (1.0 - tex_coord.y) * 2.0 - 1.0,
335 ndc_z,
336 1.0
337 );
338 return clip_to_view(view_from_clip, clip);
341/// Converts a clip-space position to a view-space position.
342/// See `bevy_light::cluster::assign::clip_to_view`.
343fn clip_to_view(view_from_clip: mat4x4<f32>, clip: vec4<f32>) -> vec4<f32> {
344 let view = view_from_clip * clip;
345 return view / view.w;
348/// Calculate the intersection of a ray from the eye through the view space
349/// position to a z plane
350/// See `bevy_light::cluster::assign::line_intersection_to_z_plane`.
351fn line_intersection_to_z_plane(origin: vec3<f32>, p: vec3<f32>, z: f32) -> vec3<f32> {
352 let v = p - origin;
353 let t = (z - dot(vec3(0.0, 0.0, 1.0), origin)) / dot(vec3(0.0, 0.0, 1.0), v);
354 return origin + t * v;
357/// Computes the near and far extents of the cluster grid.
358fn compute_z_near_and_z_far(is_orthographic: bool) -> vec2<f32> {
359 let cluster_dimensions = vec3<f32>(lights.cluster_dimensions.xyz);
360 let cluster_factors = lights.cluster_factors;
361 var z_near: f32;
362 var z_far: f32;
363 if (is_orthographic) {
364 z_near = -cluster_factors.z;
365 z_far = -(cluster_dimensions.z + cluster_factors.w * cluster_factors.z) / cluster_factors.w;
366 } else {
367 z_near = exp(cluster_factors.w / cluster_factors.z);
368 z_far = exp((cluster_dimensions.z + cluster_factors.w - 1.0) / cluster_factors.z);
369 }
370 return vec2<f32>(z_near, z_far);
373/// Returns true if a spot light should be culled.
374/// See `assign_objects_to_clusters` in `bevy_light/src/cluster/assign.rs`.
375fn cull_spot_light(
376 object_index: u32,
377 cluster_aabb_sphere_center: vec3<f32>,
378 cluster_aabb_sphere_radius: f32,
379 sphere_position: vec3<f32>,
380 sphere_radius: f32
381) -> bool {
382 let light_custom_data = clustered_lights.data[object_index].light_custom_data;
383 let light_flags = clustered_lights.data[object_index].flags;
384 let light_tan_angle = clustered_lights.data[object_index].spot_light_tan_angle;
386 // `spot_light_dir_sin_cos` in `assign_objects_to_clusters` uses
387 // `normalize(view_from_world *
388 // vec4(clusterable_object.transform.back(), 0.0).xyz)`.
389 // What we have is the XZ value of
390 // `clusterable_object.transform.forward()`. So we have to first
391 // calculate the missing Y value, then flip it to go from forward to
392 // back, then transform by the view-from-world matrix to get
393 // `view_light_direction`.
394 let world_light_direction_rev_xz = light_custom_data.xy;
395 let world_light_direction_rev_y_sign = select(
396 1.0,
397 -1.0,
398 (light_flags & POINT_LIGHT_FLAGS_SPOT_LIGHT_Y_NEGATIVE) != 0u
399 );
400 let world_light_direction_rev = vec3(
401 world_light_direction_rev_xz.x,
402 world_light_direction_rev_y_sign * sqrt(
403 1.0 -
404 world_light_direction_rev_xz.x * world_light_direction_rev_xz.x -
405 world_light_direction_rev_xz.y * world_light_direction_rev_xz.y
406 ),
407 world_light_direction_rev_xz.y
408 );
409 let world_light_direction = -world_light_direction_rev;
410 let view_light_direction = normalize((view.view_from_world *
411 vec4(world_light_direction, 0.0)).xyz);
413 let angle_cos = cos_atan(light_tan_angle);
414 let angle_sin = sin_atan(light_tan_angle);
416 // test -- based on https://bartwronski.com/2017/04/13/cull-that-cone/
417 let spot_light_offset = sphere_position - cluster_aabb_sphere_center;
418 let spot_light_dist_sq = dot(spot_light_offset, spot_light_offset);
419 let v1_len = dot(spot_light_offset, view_light_direction);
421 let distance_closest_point = angle_cos * sqrt(spot_light_dist_sq - v1_len * v1_len) -
422 v1_len * angle_sin;
423 let angle_cull = distance_closest_point > cluster_aabb_sphere_radius;
425 let front_cull = v1_len > cluster_aabb_sphere_radius + sphere_radius;
426 let back_cull = v1_len < -cluster_aabb_sphere_radius;
428 return angle_cull || front_cull || back_cull;
431/// Returns true if a rect light should be culled.
432fn cull_rect_light(
433 object_index: u32,
434 cluster_aabb_sphere_center: vec3<f32>,
435 cluster_aabb_sphere_radius: f32,
436 sphere_position: vec3<f32>
437) -> bool {
438 let rotation = clustered_lights.data[object_index].light_custom_data;
439 let world_light_direction = quat_rotate(rotation, vec3<f32>(0.0, 0.0, -1.0));
440 let view_light_direction = normalize((view.view_from_world *
441 vec4(world_light_direction, 0.0)).xyz);
443 // Signed distance from the cluster center to the light plane along the
444 // facing direction. Cull if the whole cluster sphere is on the back side.
445 let signed_distance = dot(view_light_direction, cluster_aabb_sphere_center - sphere_position);
446 return signed_distance < -cluster_aabb_sphere_radius;
449/// Computes `cos(atan(x))` cheaply.
450/// See https://en.wikipedia.org/wiki/List_of_trigonometric_identities
451fn cos_atan(tan_theta: f32) -> f32 {
452 return inverseSqrt(1.0 + tan_theta * tan_theta);
455/// Computes `sin(atan(x))` cheaply.
456/// See https://en.wikipedia.org/wiki/List_of_trigonometric_identities
457fn sin_atan(tan_theta: f32) -> f32 {
458 return tan_theta * inverseSqrt(1.0 + tan_theta * tan_theta);
461@if(!VERTEX_SHADER && POPULATE_PASS)
462/// Allocates space in the appropriate list and returns the global index that the
463/// object index should be written to.
464fn allocate_list_entry(cluster_index: u32, object_type: u32) -> u32 {
465 switch (object_type) {
466 case CLUSTERABLE_OBJECT_TYPE_POINT_LIGHT: {
467 return offsets_and_counts.data[cluster_index][0u].x +
468 atomicAdd(&scratchpad_offsets_and_counts.data[cluster_index].point_lights, 1u);
469 }
470 case CLUSTERABLE_OBJECT_TYPE_SPOT_LIGHT: {
471 return offsets_and_counts.data[cluster_index][0u].x +
472 offsets_and_counts.data[cluster_index][0u].y +
473 atomicAdd(&scratchpad_offsets_and_counts.data[cluster_index].spot_lights, 1u);
474 }
475 case CLUSTERABLE_OBJECT_TYPE_RECT_LIGHT: {
476 return offsets_and_counts.data[cluster_index][0u].x +
477 offsets_and_counts.data[cluster_index][0u].y +
478 offsets_and_counts.data[cluster_index][0u].z +
479 atomicAdd(&scratchpad_offsets_and_counts.data[cluster_index].rect_lights, 1u);
480 }
481 case CLUSTERABLE_OBJECT_TYPE_REFLECTION_PROBE: {
482 return offsets_and_counts.data[cluster_index][0u].x +
483 offsets_and_counts.data[cluster_index][0u].y +
484 offsets_and_counts.data[cluster_index][0u].z +
485 offsets_and_counts.data[cluster_index][0u].w +
486 atomicAdd(&scratchpad_offsets_and_counts.data[cluster_index].reflection_probes, 1u);
487 }
488 case CLUSTERABLE_OBJECT_TYPE_IRRADIANCE_VOLUME: {
489 return offsets_and_counts.data[cluster_index][0u].x +
490 offsets_and_counts.data[cluster_index][0u].y +
491 offsets_and_counts.data[cluster_index][0u].z +
492 offsets_and_counts.data[cluster_index][0u].w +
493 offsets_and_counts.data[cluster_index][1u].x +
494 atomicAdd(
495 &scratchpad_offsets_and_counts.data[cluster_index].irradiance_volumes,
496 1u
497 );
498 }
499 case CLUSTERABLE_OBJECT_TYPE_DECAL: {
500 return offsets_and_counts.data[cluster_index][0u].x +
501 offsets_and_counts.data[cluster_index][0u].y +
502 offsets_and_counts.data[cluster_index][0u].z +
503 offsets_and_counts.data[cluster_index][0u].w +
504 offsets_and_counts.data[cluster_index][1u].x +
505 offsets_and_counts.data[cluster_index][1u].y +
506 atomicAdd(&scratchpad_offsets_and_counts.data[cluster_index].decals, 1u);
507 }
508 default: {}
509 }
510 return 0xffffffffu;
512@if(!VERTEX_SHADER && !POPULATE_PASS)
513/// Increments the count of objects of the given type for the given cluster.
514fn increment_object_count(cluster_index: u32, object_type: u32) {
515 switch (object_type) {
516 case CLUSTERABLE_OBJECT_TYPE_POINT_LIGHT: {
517 atomicAdd(&offsets_and_counts.data[cluster_index].point_lights, 1u);
518 }
519 case CLUSTERABLE_OBJECT_TYPE_SPOT_LIGHT: {
520 atomicAdd(&offsets_and_counts.data[cluster_index].spot_lights, 1u);
521 }
522 case CLUSTERABLE_OBJECT_TYPE_RECT_LIGHT: {
523 atomicAdd(&offsets_and_counts.data[cluster_index].rect_lights, 1u);
524 }
525 case CLUSTERABLE_OBJECT_TYPE_REFLECTION_PROBE: {
526 atomicAdd(&offsets_and_counts.data[cluster_index].reflection_probes, 1u);
527 }
528 case CLUSTERABLE_OBJECT_TYPE_IRRADIANCE_VOLUME: {
529 atomicAdd(&offsets_and_counts.data[cluster_index].irradiance_volumes, 1u);
530 }
531 case CLUSTERABLE_OBJECT_TYPE_DECAL: {
532 atomicAdd(&offsets_and_counts.data[cluster_index].decals, 1u);
533 }
534 default: {}
535 }
538/// Looks up and returns the world-space center and radius of the bounding sphere
539/// for the object with the given index and type.
540/// Returns a 4-vector with the fields `vec4(center X, center Y, center Z, radius)`.
541fn get_object_bounding_sphere(object_index: u32, object_type: u32) -> vec4<f32> {
542 var position = vec3<f32>(0.0);
543 var radius = 0.0;
544 switch (object_type) {
545 case CLUSTERABLE_OBJECT_TYPE_POINT_LIGHT, CLUSTERABLE_OBJECT_TYPE_SPOT_LIGHT, CLUSTERABLE_OBJECT_TYPE_RECT_LIGHT: {
546 position = clustered_lights.data[object_index].position_radius.xyz;
547 radius = clustered_lights.data[object_index].range;
548 }
549 case CLUSTERABLE_OBJECT_TYPE_REFLECTION_PROBE: {
550 position = light_probes.reflection_probes[object_index].world_position;
551 radius = light_probes.reflection_probes[object_index].bounding_sphere_radius;
552 }
553 case CLUSTERABLE_OBJECT_TYPE_IRRADIANCE_VOLUME: {
554 position = light_probes.irradiance_volumes[object_index].world_position;
555 radius = light_probes.irradiance_volumes[object_index].bounding_sphere_radius;
556 }
557 case CLUSTERABLE_OBJECT_TYPE_DECAL: {
558 position = clustered_decals.decals[object_index].world_position;
559 radius = clustered_decals.decals[object_index].bounding_sphere_radius;
560 }
561 default: {}
562 }
563 return vec4(position, radius);
Contents