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(),
         )),
     );
@@ -317,7 +311,9 @@ pub fn solari_lighting(
 
     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 + 4,
+        ),
         "solari_lighting/world_cache_active_cells_count",
     );
 }
@@ -347,15 +343,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,48 @@ 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);
 
+/// Layout constants for the packed `WorldCache` buffer.
+///
+/// The `ShaderType` mirror lives in this module and is intentionally private so it cannot be
+/// constructed, because it would be too large.
+mod world_cache_layout {
+    use bevy_math::{Vec3, Vec4};
+    use bevy_render::render_resource::ShaderType;
+
+    const WORLD_CACHE_LEN: usize = super::WORLD_CACHE_SIZE as usize;
+
+    #[derive(ShaderType)]
+    struct GeometryData {
+        world_position: Vec3,
+        padding_a: u32,
+        world_normal: Vec3,
+        padding_b: u32,
+    }
+
+    #[derive(ShaderType)]
+    struct WorldCache {
+        checksums: [u32; WORLD_CACHE_LEN],
+        life: [u32; WORLD_CACHE_LEN],
+        radiance: [Vec4; WORLD_CACHE_LEN],
+        geometry_data: [GeometryData; WORLD_CACHE_LEN],
+        luminance_deltas: [f32; WORLD_CACHE_LEN],
+        active_cells_new_radiance: [Vec3; WORLD_CACHE_LEN],
+        a: [u32; WORLD_CACHE_LEN],
+        b: [u32; WORLD_CACHE_LEN / 1024],
+        active_cell_indices: [u32; WORLD_CACHE_LEN],
+        active_cells_count: u32,
+    }
+
+    // `ShaderType::METADATA` is internal, but we need the field offset and size without
+    // constructing this large layout type.
+    pub const ACTIVE_CELLS_COUNT_OFFSET: u64 = WorldCache::METADATA.last_offset();
+    /// Must stay under wgpu's default `max_storage_buffer_binding_size` (128 MiB or 2^27 bytes).
+    pub const BUFFER_SIZE: u64 = WorldCache::METADATA.min_size().get();
+}
+
+pub const WORLD_CACHE_ACTIVE_CELLS_COUNT_OFFSET: u64 =
+    world_cache_layout::ACTIVE_CELLS_COUNT_OFFSET;
+
 /// GPU representation of the user-configurable [`SolariLighting`] settings, plus
 /// per-frame state.
 ///
@@ -93,16 +135,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 +229,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_layout::BUFFER_SIZE,
             usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
             mapped_at_creation: false,
         });
@@ -279,16 +249,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

@@ -19,21 +19,25 @@ enable wgpu_ray_query;
 @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}>;
+// Keep in sync with `world_cache_layout::WorldCache` in `prepare.rs`.
+struct WorldCache {
+    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}>;
+    life: array<u32, #{WORLD_CACHE_SIZE}>,
 #else
-@group(1) @binding(13) var<storage, read_write> world_cache_life: array<atomic<u32>, #{WORLD_CACHE_SIZE}>;
+    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;
+    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,
+}
+@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>;

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