Solari support on metal backends

_release-content/release-notes/solari_metal.md

@@ -0,0 +1,15 @@
+---
+title: Solari on Metal
+authors: ["@mate-h"]
+pull_requests: [25123]
+---
+
+Ray tracing on Metal has been available since wgpu 29, and with bindless storage buffers landing in wgpu 30, Solari now runs on Apple Silicon Macs.
+
+Run the Solari example on a compatible Mac:
+
+```sh
+cargo run --example solari --features bevy_solari,https,free_camera
+```
+
+Denoising is not available on Metal yet, DLSS is NVIDIA-only. MetalFX Ray Reconstruction or Open Image Denoise 3.0 are promising paths for cross-platform denoising in the future.

crates/bevy_solari/src/realtime/node.rs

@@ -1,4 +1,7 @@
-use super::prepare::{SolariLightingResources, LIGHT_TILE_BLOCKS, WORLD_CACHE_SIZE};
+use super::prepare::{
+    SolariLightingResources, LIGHT_TILE_BLOCKS, WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET,
+    WORLD_CACHE_SIZE,
+};
 use crate::scene::RaytracingSceneBindings;
 #[cfg(all(feature = "dlss", not(feature = "force_disable_dlss")))]
 use bevy_anti_alias::dlss::ViewDlssRayReconstructionTextures;
@@ -179,16 +182,7 @@ pub fn solari_lighting(
             previous_depth_buffer,
             view_uniforms_binding,
             previous_view_uniforms_binding,
-            s.world_cache_checksums.as_entire_binding(),
-            s.world_cache_life.as_entire_binding(),
-            s.world_cache_radiance.as_entire_binding(),
-            s.world_cache_geometry_data.as_entire_binding(),
-            s.world_cache_luminance_deltas.as_entire_binding(),
-            s.world_cache_active_cells_new_radiance.as_entire_binding(),
-            s.world_cache_a.as_entire_binding(),
-            s.world_cache_b.as_entire_binding(),
-            s.world_cache_active_cell_indices.as_entire_binding(),
-            s.world_cache_active_cells_count.as_entire_binding(),
+            s.world_cache.as_entire_binding(),
             s.constants.as_entire_binding(),
         )),
     );
@@ -315,9 +309,13 @@ pub fn solari_lighting(
 
     drop(pass);
 
+    // Active cell count readback.
     diagnostics.record_u32(
         ctx.command_encoder(),
-        &s.world_cache_active_cells_count.slice(..),
+        &s.world_cache.slice(
+            WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET
+                ..WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET + size_of::<u32>() as u64,
+        ),
         "solari_lighting/world_cache_active_cells_count",
     );
 }
@@ -347,15 +345,6 @@ pub fn init_solari_lighting_pipelines(
                 uniform_buffer::<ViewUniform>(true),
                 uniform_buffer::<PreviousViewData>(true),
                 storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
-                storage_buffer_sized(false, None),
                 uniform_buffer_sized(false, None),
             ),
         ),

crates/bevy_solari/src/realtime/prepare.rs

@@ -44,6 +44,20 @@ pub const LIGHT_TILE_SAMPLES_PER_BLOCK: u64 = 1024;
 /// Amount of entries in the world cache (must be a power of 2, and >= 2^10)
 pub const WORLD_CACHE_SIZE: u64 = 2u64.pow(20);
 
+/// Sum of per-cell field sizes in `WorldCache`. Keep in sync with `realtime_bindings.wgsl`.
+const WORLD_CACHE_ENTRY_SIZE: u64 = 84;
+
+/// Size of the fixed `b` array (`array<u32, WORLD_CACHE_SIZE / 1024>`).
+const WORLD_CACHE_B_SIZE: u64 = (WORLD_CACHE_SIZE / 1024) * size_of::<u32>() as u64;
+
+/// Offset of `active_cells_count`.
+pub const WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET: u64 =
+    WORLD_CACHE_SIZE * WORLD_CACHE_ENTRY_SIZE + WORLD_CACHE_B_SIZE;
+
+/// Must stay under wgpu's default `max_storage_buffer_binding_size` (128 MiB or 2^27 bytes).
+pub const WORLD_CACHE_BUFFER_SIZE: u64 =
+    WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET + size_of::<u32>() as u64;
+
 /// GPU representation of the user-configurable [`SolariLighting`] settings, plus
 /// per-frame state.
 ///
@@ -93,16 +107,7 @@ pub struct SolariLightingResources {
     pub light_tile_resolved_samples: Buffer,
     pub reservoirs_a: Buffer,
     pub reservoirs_b: Buffer,
-    pub world_cache_checksums: Buffer,
-    pub world_cache_life: Buffer,
-    pub world_cache_radiance: Buffer,
-    pub world_cache_geometry_data: Buffer,
-    pub world_cache_luminance_deltas: Buffer,
-    pub world_cache_active_cells_new_radiance: Buffer,
-    pub world_cache_a: Buffer,
-    pub world_cache_b: Buffer,
-    pub world_cache_active_cell_indices: Buffer,
-    pub world_cache_active_cells_count: Buffer,
+    pub world_cache: Buffer,
     pub world_cache_active_cells_dispatch: Buffer,
     pub view_size: UVec2,
 }
@@ -196,72 +201,9 @@ pub fn prepare_solari_lighting_resources(
         let reservoirs_a = reservoirs_buffer("solari_lighting_reservoirs_a");
         let reservoirs_b = reservoirs_buffer("solari_lighting_reservoirs_b");
 
-        let world_cache_checksums = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_checksums"),
-            size: WORLD_CACHE_SIZE * size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_life = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_life"),
-            size: WORLD_CACHE_SIZE * size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_radiance = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_radiance"),
-            size: WORLD_CACHE_SIZE * size_of::<[f32; 4]>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_geometry_data = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_geometry_data"),
-            size: WORLD_CACHE_SIZE * size_of::<[f32; 8]>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_luminance_deltas = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_luminance_deltas"),
-            size: WORLD_CACHE_SIZE * size_of::<f32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_active_cells_new_radiance =
-            render_device.create_buffer(&BufferDescriptor {
-                label: Some("solari_lighting_world_cache_active_cells_new_radiance"),
-                size: WORLD_CACHE_SIZE * size_of::<[f32; 4]>() as u64,
-                usage: BufferUsages::STORAGE,
-                mapped_at_creation: false,
-            });
-
-        let world_cache_a = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_a"),
-            size: WORLD_CACHE_SIZE * size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-        let world_cache_b = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_b"),
-            size: 1024 * size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_active_cell_indices = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_active_cell_indices"),
-            size: WORLD_CACHE_SIZE * size_of::<u32>() as u64,
-            usage: BufferUsages::STORAGE,
-            mapped_at_creation: false,
-        });
-
-        let world_cache_active_cells_count = render_device.create_buffer(&BufferDescriptor {
-            label: Some("solari_lighting_world_cache_active_cells_count"),
-            size: size_of::<u32>() as u64,
+        let world_cache = render_device.create_buffer(&BufferDescriptor {
+            label: Some("solari_lighting_world_cache"),
+            size: WORLD_CACHE_BUFFER_SIZE,
             usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
             mapped_at_creation: false,
         });
@@ -279,16 +221,7 @@ pub fn prepare_solari_lighting_resources(
             light_tile_resolved_samples,
             reservoirs_a,
             reservoirs_b,
-            world_cache_checksums,
-            world_cache_life,
-            world_cache_radiance,
-            world_cache_geometry_data,
-            world_cache_luminance_deltas,
-            world_cache_active_cells_new_radiance,
-            world_cache_a,
-            world_cache_b,
-            world_cache_active_cell_indices,
-            world_cache_active_cells_count,
+            world_cache,
             world_cache_active_cells_dispatch,
             view_size,
         });

crates/bevy_solari/src/realtime/realtime_bindings.wgsl

@@ -18,22 +18,8 @@ enable wgpu_ray_query;
 @group(1) @binding(9) var previous_depth_buffer: texture_depth_2d;
 @group(1) @binding(10) var<uniform> view: View;
 @group(1) @binding(11) var<uniform> previous_view: PreviousViewUniforms;
-
-@group(1) @binding(12) var<storage, read_write> world_cache_checksums: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
-#ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER
-@group(1) @binding(13) var<storage, read_write> world_cache_life: array<u32, #{WORLD_CACHE_SIZE}>;
-#else
-@group(1) @binding(13) var<storage, read_write> world_cache_life: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
-#endif
-@group(1) @binding(14) var<storage, read_write> world_cache_radiance: array<vec4<f32>, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(15) var<storage, read_write> world_cache_geometry_data: array<WorldCacheGeometryData, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(16) var<storage, read_write> world_cache_luminance_deltas: array<f32, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(17) var<storage, read_write> world_cache_active_cells_new_radiance: array<vec3<f32>, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(18) var<storage, read_write> world_cache_a: array<u32, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(19) var<storage, read_write> world_cache_b: array<u32, 1024u>;
-@group(1) @binding(20) var<storage, read_write> world_cache_active_cell_indices: array<u32, #{WORLD_CACHE_SIZE}>;
-@group(1) @binding(21) var<storage, read_write> world_cache_active_cells_count: u32;
-@group(1) @binding(22) var<uniform> constants: SolariLightingSettings;
+@group(1) @binding(12) var<storage, read_write> world_cache: WorldCache;
+@group(1) @binding(13) var<uniform> constants: SolariLightingSettings;
 
 #ifdef DLSS_RR_GUIDE_BUFFERS
 @group(2) @binding(0) var diffuse_albedo: texture_storage_2d<rgba8unorm, write>;
@@ -96,3 +82,21 @@ struct WorldCacheGeometryData {
     world_normal: vec3<f32>,
     padding_b: u32,
 }
+
+// Keep in sync with the packed world-cache layout constants in `prepare.rs`.
+struct WorldCache {
+    checksums: array<atomic<u32>, #{WORLD_CACHE_SIZE}>,
+#ifdef WORLD_CACHE_NON_ATOMIC_LIFE_BUFFER
+    life: array<u32, #{WORLD_CACHE_SIZE}>,
+#else
+    life: array<atomic<u32>, #{WORLD_CACHE_SIZE}>,
+#endif
+    radiance: array<vec4<f32>, #{WORLD_CACHE_SIZE}>,
+    geometry_data: array<WorldCacheGeometryData, #{WORLD_CACHE_SIZE}>,
+    luminance_deltas: array<f32, #{WORLD_CACHE_SIZE}>,
+    active_cells_new_radiance: array<vec3<f32>, #{WORLD_CACHE_SIZE}>,
+    a: array<u32, #{WORLD_CACHE_SIZE}>,
+    b: array<u32, 1024u>,
+    active_cell_indices: array<u32, #{WORLD_CACHE_SIZE}>,
+    active_cells_count: u32,
+}

crates/bevy_solari/src/realtime/world_cache_compact.wgsl

@@ -1,16 +1,7 @@
 enable wgpu_ray_query;
 
 #import bevy_solari::world_cache::WORLD_CACHE_EMPTY_CELL
-#import bevy_solari::realtime_bindings::{
-    world_cache_life,
-    world_cache_checksums,
-    world_cache_radiance,
-    world_cache_luminance_deltas,
-    world_cache_a,
-    world_cache_b,
-    world_cache_active_cell_indices,
-    world_cache_active_cells_count,
-}
+#import bevy_solari::realtime_bindings::world_cache
 
 @group(2) @binding(0) var<storage, read_write> world_cache_active_cells_dispatch: vec3<u32>;
 
@@ -19,15 +10,15 @@ var<workgroup> w2: array<u32, 1024u>;
 
 @compute @workgroup_size(1024, 1, 1)
 fn decay_world_cache(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    var life = world_cache_life[global_id.x];
+    var life = world_cache.life[global_id.x];
     if life > 0u {
         life -= 1u;
-        world_cache_life[global_id.x] = life;
+        world_cache.life[global_id.x] = life;
 
         if life == 0u {
-            world_cache_checksums[global_id.x] = WORLD_CACHE_EMPTY_CELL;
-            world_cache_radiance[global_id.x] = vec4(0.0);
-            world_cache_luminance_deltas[global_id.x] = 0.0;
+            world_cache.checksums[global_id.x] = WORLD_CACHE_EMPTY_CELL;
+            world_cache.radiance[global_id.x] = vec4(0.0);
+            world_cache.luminance_deltas[global_id.x] = 0.0;
         }
     }
 }
@@ -37,7 +28,7 @@ fn compact_world_cache_single_block(
     @builtin(global_invocation_id) cell_id: vec3<u32>,
     @builtin(local_invocation_index) t: u32,
 ) {
-    if t == 0u { w1[0u] = 0u; } else { w1[t] = u32(world_cache_life[cell_id.x - 1u] != 0u); }; workgroupBarrier();
+    if t == 0u { w1[0u] = 0u; } else { w1[t] = u32(world_cache.life[cell_id.x - 1u] != 0u); }; workgroupBarrier();
     if t < 1u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 1u]; } workgroupBarrier();
     if t < 2u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 2u]; } workgroupBarrier();
     if t < 4u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 4u]; } workgroupBarrier();
@@ -47,12 +38,12 @@ fn compact_world_cache_single_block(
     if t < 64u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 64u]; } workgroupBarrier();
     if t < 128u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 128u]; } workgroupBarrier();
     if t < 256u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 256u]; } workgroupBarrier();
-    if t < 512u { world_cache_a[cell_id.x] = w2[t]; } else { world_cache_a[cell_id.x] = w2[t] + w2[t - 512u]; }
+    if t < 512u { world_cache.a[cell_id.x] = w2[t]; } else { world_cache.a[cell_id.x] = w2[t] + w2[t - 512u]; }
 }
 
 @compute @workgroup_size(1024, 1, 1)
 fn compact_world_cache_blocks(@builtin(local_invocation_index) t: u32) {
-    if t == 0u { w1[0u] = 0u; } else { w1[t] = world_cache_a[t * 1024u - 1u]; }; workgroupBarrier();
+    if t == 0u { w1[0u] = 0u; } else { w1[t] = world_cache.a[t * 1024u - 1u]; }; workgroupBarrier();
     if t < 1u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 1u]; } workgroupBarrier();
     if t < 2u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 2u]; } workgroupBarrier();
     if t < 4u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 4u]; } workgroupBarrier();
@@ -62,7 +53,7 @@ fn compact_world_cache_blocks(@builtin(local_invocation_index) t: u32) {
     if t < 64u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 64u]; } workgroupBarrier();
     if t < 128u { w1[t] = w2[t]; } else { w1[t] = w2[t] + w2[t - 128u]; } workgroupBarrier();
     if t < 256u { w2[t] = w1[t]; } else { w2[t] = w1[t] + w1[t - 256u]; } workgroupBarrier();
-    if t < 512u { world_cache_b[t] = w2[t]; } else { world_cache_b[t] = w2[t] + w2[t - 512u]; }
+    if t < 512u { world_cache.b[t] = w2[t]; } else { world_cache.b[t] = w2[t] + w2[t - 512u]; }
 }
 
 @compute @workgroup_size(1024, 1, 1)
@@ -71,16 +62,16 @@ fn compact_world_cache_write_active_cells(
     @builtin(workgroup_id) workgroup_id: vec3<u32>,
     @builtin(local_invocation_index) thread_index: u32,
 ) {
-    let compacted_index = world_cache_a[cell_id.x] + world_cache_b[workgroup_id.x];
-    let cell_active = world_cache_life[cell_id.x] != 0u;
+    let compacted_index = world_cache.a[cell_id.x] + world_cache.b[workgroup_id.x];
+    let cell_active = world_cache.life[cell_id.x] != 0u;
 
     if cell_active {
-        world_cache_active_cell_indices[compacted_index] = cell_id.x;
+        world_cache.active_cell_indices[compacted_index] = cell_id.x;
     }
 
     if thread_index == 1023u && workgroup_id.x == 1023u {
         let active_cell_count = compacted_index + u32(cell_active);
-        world_cache_active_cells_count = active_cell_count;
+        world_cache.active_cells_count = active_cell_count;
         world_cache_active_cells_dispatch = vec3((active_cell_count + 63u) / 64u, 1u, 1u);
     }
 }

crates/bevy_solari/src/realtime/world_cache_query.wgsl

@@ -4,13 +4,7 @@ enable wgpu_ray_query;
 
 #import bevy_pbr::utils::{rand_f, rand_vec2f}
 #import bevy_render::maths::orthonormalize
-#import bevy_solari::realtime_bindings::{
-    world_cache_life,
-    world_cache_checksums,
-    world_cache_radiance,
-    world_cache_geometry_data,
-    constants,
-}
+#import bevy_solari::realtime_bindings::{world_cache, constants}
 
 /// Maximum amount of frames a cell can live for without being queried
 const WORLD_CACHE_CELL_LIFETIME: u32 = 10u;
@@ -40,25 +34,25 @@ fn query_world_cache(world_position_in: vec3<f32>, world_normal: vec3<f32>, view
     let checksum = compute_checksum(world_position_quantized, world_normal_quantized);
 
     for (var i = 0u; i < WORLD_CACHE_MAX_SEARCH_STEPS; i++) {
-        let cas = atomicCompareExchangeWeak(&world_cache_checksums[key], WORLD_CACHE_EMPTY_CELL, checksum);
+        let cas = atomicCompareExchangeWeak(&world_cache.checksums[key], WORLD_CACHE_EMPTY_CELL, checksum);
         let existing_checksum = cas.old_value;
 
         // Cell already exists or is empty - reset lifetime
         if existing_checksum == checksum || existing_checksum == WORLD_CACHE_EMPTY_CELL {
 #ifndef WORLD_CACHE_QUERY_ATOMIC_MAX_LIFETIME
-            atomicStore(&world_cache_life[key], cell_lifetime);
+            atomicStore(&world_cache.life[key], cell_lifetime);
 #else
-            atomicMax(&world_cache_life[key], cell_lifetime);
+            atomicMax(&world_cache.life[key], cell_lifetime);
 #endif
         }
 
         if existing_checksum == checksum {
             // Cache entry already exists - get radiance
-            return world_cache_radiance[key].rgb;
+            return world_cache.radiance[key].rgb;
         } else if existing_checksum == WORLD_CACHE_EMPTY_CELL && cas.exchanged {
             // Cell is empty - initialize it
-            world_cache_geometry_data[key].world_position = world_position;
-            world_cache_geometry_data[key].world_normal = world_normal;
+            world_cache.geometry_data[key].world_position = world_position;
+            world_cache.geometry_data[key].world_normal = world_normal;
             return vec3(0.0);
         } else {
             // Collision - linear probe to next entry