Merge c589cd2b39 into 877d278785

crates/bevy_light/src/lib.rs

@@ -27,7 +27,7 @@ use cluster::{
 mod ambient_light;
 pub use ambient_light::AmbientLight;
 mod probe;
-pub use probe::{EnvironmentMapLight, IrradianceVolume, LightProbe};
+pub use probe::{EnvironmentMapLight, GeneratedEnvironmentMapLight, IrradianceVolume, LightProbe};
 mod volumetric;
 pub use volumetric::{FogVolume, VolumetricFog, VolumetricLight};
 pub mod cascade;

crates/bevy_light/src/probe.rs

@@ -108,6 +108,47 @@ impl Default for EnvironmentMapLight {
     }
 }
 
+/// A generated environment map that is filtered at runtime.
+///
+/// See `bevy_pbr::light_probe::generate` for detailed information.
+#[derive(Clone, Component, Reflect)]
+#[reflect(Component, Default, Clone)]
+pub struct GeneratedEnvironmentMapLight {
+    /// Source cubemap to be filtered on the GPU, size must be a power of two.
+    pub environment_map: Handle<Image>,
+
+    /// Scale factor applied to the diffuse and specular light generated by this
+    /// component. Expressed in cd/m² (candela per square meter).
+    pub intensity: f32,
+
+    /// World-space rotation applied to the cubemap.
+    pub rotation: Quat,
+
+    /// Whether this light contributes diffuse lighting to meshes that already
+    /// have baked lightmaps.
+    pub affects_lightmapped_mesh_diffuse: bool,
+
+    /// White point applied during reverse tone-mapping.
+    ///
+    /// This value attenuates extremely bright texels in the source cubemap to
+    /// suppress *fireflies* introduced by HDR highlights. Increasing
+    /// `white_point` preserves more highlight energy; lowering it clamps
+    /// highlights more aggressively.
+    pub white_point: f32,
+}
+
+impl Default for GeneratedEnvironmentMapLight {
+    fn default() -> Self {
+        GeneratedEnvironmentMapLight {
+            environment_map: Handle::default(),
+            intensity: 0.0,
+            rotation: Quat::IDENTITY,
+            affects_lightmapped_mesh_diffuse: true,
+            white_point: 1.0,
+        }
+    }
+}
+
 /// The component that defines an irradiance volume.
 ///
 /// See `bevy_pbr::irradiance_volume` for detailed information.

crates/bevy_pbr/src/lib.rs

@@ -88,8 +88,8 @@ pub mod prelude {
     };
     #[doc(hidden)]
     pub use bevy_light::{
-        light_consts, AmbientLight, DirectionalLight, EnvironmentMapLight, LightProbe, PointLight,
-        SpotLight,
+        light_consts, AmbientLight, DirectionalLight, EnvironmentMapLight,
+        GeneratedEnvironmentMapLight, LightProbe, PointLight, SpotLight,
     };
 }
 

crates/bevy_pbr/src/light_probe/copy_mip0.wgsl

@@ -0,0 +1,21 @@
+// Copy the base mip (level 0) from a source cubemap to a destination cubemap,
+// performing format conversion if needed (the destination is always rgba16float).
+// The alpha channel is filled with 1.0.
+
+@group(0) @binding(0) var src_cubemap: texture_2d_array<f32>;
+@group(0) @binding(1) var dst_cubemap: texture_storage_2d_array<rgba16float, write>;
+
+@compute
+@workgroup_size(8, 8, 1)
+fn copy_mip0(@builtin(global_invocation_id) gid: vec3u) {
+    let size = textureDimensions(src_cubemap).xy;
+
+    // Bounds check
+    if (any(gid.xy >= size)) {
+        return;
+    }
+
+    let color = textureLoad(src_cubemap, vec2u(gid.xy), gid.z, 0);
+
+    textureStore(dst_cubemap, vec2u(gid.xy), gid.z, vec4f(color.rgb, 1.0));
+} 

crates/bevy_pbr/src/light_probe/environment_filter.wgsl

@@ -0,0 +1,341 @@
+#import bevy_render::maths::{PI, PI_2};
+#import bevy_pbr::utils::{build_orthonormal_basis};
+
+struct FilteringConstants {
+    mip_level: f32,
+    sample_count: u32,
+    roughness: f32,
+    blue_noise_size: vec2f,
+    white_point: f32,
+}
+
+@group(0) @binding(0) var input_texture: texture_2d_array<f32>;
+@group(0) @binding(1) var input_sampler: sampler;
+@group(0) @binding(2) var output_texture: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(3) var<uniform> constants: FilteringConstants;
+@group(0) @binding(4) var blue_noise_texture: texture_2d<f32>;
+
+// Tonemapping functions to reduce fireflies
+fn tonemap(color: vec3f) -> vec3f {
+    return color / (color + vec3(constants.white_point));
+}
+fn reverse_tonemap(color: vec3f) -> vec3f {
+    return constants.white_point * color / (vec3(1.0) - color);
+}
+
+// Convert UV and face index to direction vector
+fn sample_cube_dir(uv: vec2f, face: u32) -> vec3f {
+    // Convert from [0,1] to [-1,1]
+    let uvc = 2.0 * uv - 1.0;
+    
+    // Generate direction based on the cube face
+    var dir: vec3f;
+    switch(face) {
+        case 0u: { dir = vec3f( 1.0,  -uvc.y, -uvc.x); } // +X
+        case 1u: { dir = vec3f(-1.0,  -uvc.y,  uvc.x); } // -X
+        case 2u: { dir = vec3f( uvc.x,  1.0,   uvc.y); } // +Y
+        case 3u: { dir = vec3f( uvc.x, -1.0,  -uvc.y); } // -Y
+        case 4u: { dir = vec3f( uvc.x, -uvc.y,  1.0);  } // +Z
+        case 5u: { dir = vec3f(-uvc.x, -uvc.y, -1.0);  } // -Z
+        default: { dir = vec3f(0.0); }
+    }
+    return normalize(dir);
+}
+
+// Convert direction vector to cube face UV
+struct CubeUV {
+    uv: vec2f,
+    face: u32,
+}
+fn dir_to_cube_uv(dir: vec3f) -> CubeUV {
+    let abs_dir = abs(dir);
+    var face: u32 = 0u;
+    var uv: vec2f = vec2f(0.0);
+    
+    // Find the dominant axis to determine face
+    if (abs_dir.x >= abs_dir.y && abs_dir.x >= abs_dir.z) {
+        // X axis is dominant
+        if (dir.x > 0.0) {
+            face = 0u; // +X
+            uv = vec2f(-dir.z, -dir.y) / dir.x;
+        } else {
+            face = 1u; // -X
+            uv = vec2f(dir.z, -dir.y) / abs_dir.x;
+        }
+    } else if (abs_dir.y >= abs_dir.x && abs_dir.y >= abs_dir.z) {
+        // Y axis is dominant
+        if (dir.y > 0.0) {
+            face = 2u; // +Y
+            uv = vec2f(dir.x, dir.z) / dir.y;
+        } else {
+            face = 3u; // -Y
+            uv = vec2f(dir.x, -dir.z) / abs_dir.y;
+        }
+    } else {
+        // Z axis is dominant
+        if (dir.z > 0.0) {
+            face = 4u; // +Z
+            uv = vec2f(dir.x, -dir.y) / dir.z;
+        } else {
+            face = 5u; // -Z
+            uv = vec2f(-dir.x, -dir.y) / abs_dir.z;
+        }
+    }
+    
+    // Convert from [-1,1] to [0,1]
+    return CubeUV(uv * 0.5 + 0.5, face);
+}
+
+// Sample an environment map with a specific LOD
+fn sample_environment(dir: vec3f, level: f32) -> vec4f {
+    let cube_uv = dir_to_cube_uv(dir);
+    return textureSampleLevel(input_texture, input_sampler, cube_uv.uv, cube_uv.face, level);
+}
+
+// Hammersley sequence for quasi-random points
+fn hammersley_2d(i: u32, n: u32) -> vec2f {
+    let inv_n = 1.0 / f32(n);
+    let vdc = f32(reverseBits(i)) * 2.3283064365386963e-10; // 1/2^32
+    return vec2f(f32(i) * inv_n, vdc);
+}
+
+// Blue noise randomization
+fn sample_noise(pixel_coords: vec2u) -> vec4f {
+    let noise_size = vec2u(u32(constants.blue_noise_size.x), u32(constants.blue_noise_size.y));
+    let noise_coords = pixel_coords % noise_size;
+    let uv = vec2f(noise_coords) / constants.blue_noise_size;
+    return textureSampleLevel(blue_noise_texture, input_sampler, uv, 0.0);
+}
+
+// from bevy_pbr/src/render/pbr_lighting.wgsl
+fn perceptualRoughnessToRoughness(perceptualRoughness: f32) -> f32 {
+    // clamp perceptual roughness to prevent precision problems
+    // According to Filament design 0.089 is recommended for mobile
+    // Filament uses 0.045 for non-mobile
+    let clampedPerceptualRoughness = clamp(perceptualRoughness, 0.089, 1.0);
+    return clampedPerceptualRoughness * clampedPerceptualRoughness;
+}
+
+// GGX/Trowbridge-Reitz normal distribution function (D term)
+// from bevy_pbr/src/render/pbr_lighting.wgsl
+fn D_GGX(roughness: f32, NdotH: f32) -> f32 {
+    let oneMinusNdotHSquared = 1.0 - NdotH * NdotH;
+    let a = NdotH * roughness;
+    let k = roughness / (oneMinusNdotHSquared + a * a);
+    let d = k * k * (1.0 / PI);
+    return d;
+}
+
+// Importance sample GGX normal distribution function for a given roughness
+fn importance_sample_ggx(xi: vec2f, roughness: f32, normal: vec3f) -> vec3f {
+    // Sample in spherical coordinates
+    let phi = PI_2 * xi.x;
+    
+    // GGX mapping from uniform random to GGX distribution
+    let cos_theta = sqrt((1.0 - xi.y) / (1.0 + (roughness * roughness - 1.0) * xi.y));
+    let sin_theta = sqrt(1.0 - cos_theta * cos_theta);
+    
+    // Convert to cartesian
+    let h = vec3f(
+        sin_theta * cos(phi),
+        sin_theta * sin(phi),
+        cos_theta
+    );
+    
+    // Transform from tangent to world space
+    return build_orthonormal_basis(normal) * h;
+}
+
+// Calculate LOD for environment map lookup using filtered importance sampling
+fn calculate_environment_map_lod(pdf: f32, width: f32, samples: f32) -> f32 {
+    // Solid angle of current sample
+    let omega_s = 1.0 / (samples * pdf);
+    
+    // Solid angle of a texel in the environment map
+    let omega_p = 4.0 * PI / (6.0 * width * width);
+    
+    // Filtered importance sampling: compute the correct LOD
+    return 0.5 * log2(omega_s / omega_p);
+}
+
+// Smith geometric shadowing function
+fn G_Smith(NoV: f32, NoL: f32, roughness: f32) -> f32 {
+    let k = roughness / 2.0;
+    let GGXL = NoL / (NoL * (1.0 - k) + k);
+    let GGXV = NoV / (NoV * (1.0 - k) + k);
+    return GGXL * GGXV;
+}
+
+@compute
+@workgroup_size(8, 8, 1)
+fn generate_radiance_map(@builtin(global_invocation_id) global_id: vec3u) {
+    let size = textureDimensions(output_texture).xy;
+    let invSize = 1.0 / vec2f(size);
+    
+    let coords = vec2u(global_id.xy);
+    let face = global_id.z;
+    
+    if (any(coords >= size)) {
+        return;
+    }
+    
+    // Convert texture coordinates to direction vector
+    let uv = (vec2f(coords) + 0.5) * invSize;
+    let normal = sample_cube_dir(uv, face);
+    
+    // For radiance map, view direction = normal for perfect reflection
+    let view = normal;
+    
+    // Convert perceptual roughness to physical microfacet roughness
+    let perceptual_roughness = constants.roughness;
+    let roughness = perceptualRoughnessToRoughness(perceptual_roughness);
+    
+    // Get blue noise offset for stratification
+    let vector_noise = sample_noise(coords);
+    
+    var radiance = vec3f(0.0);
+    var total_weight = 0.0;
+    
+    // Skip sampling for mirror reflection (roughness = 0)
+    if (roughness < 0.01) {
+        radiance = sample_environment(normal, 0.0).rgb;
+        textureStore(output_texture, coords, face, vec4f(radiance, 1.0));
+        return;
+    }
+    
+    // For higher roughness values, use importance sampling
+    let sample_count = constants.sample_count;
+    
+    for (var i = 0u; i < sample_count; i++) {
+        // Get sample coordinates from Hammersley sequence with blue noise offset
+        var xi = hammersley_2d(i, sample_count);
+        xi = fract(xi + vector_noise.rg); // Apply Cranley-Patterson rotation
+        
+        // Sample the GGX distribution to get a half vector
+        let half_vector = importance_sample_ggx(xi, roughness, normal);
+        
+        // Calculate reflection vector from half vector
+        let light_dir = reflect(-view, half_vector);
+        
+        // Calculate weight (N·L)
+        let NoL = dot(normal, light_dir);
+        
+        if (NoL > 0.0) {
+            // Calculate values needed for PDF
+            let NoH = dot(normal, half_vector);
+            let VoH = dot(view, half_vector);
+            let NoV = dot(normal, view);
+            
+            // Get the geometric shadowing term
+            let G = G_Smith(NoV, NoL, roughness);
+            
+            // Probability Distribution Function
+            let pdf = D_GGX(roughness, NoH) * NoH / (4.0 * VoH);
+            
+            // Calculate LOD using filtered importance sampling
+            // This is crucial to avoid fireflies and improve quality
+            let width = f32(size.x);
+            let lod = calculate_environment_map_lod(pdf, width, f32(sample_count));
+            
+            // Get source mip level - ensure we don't go negative
+            let source_mip = max(0.0, lod);
+            
+            // Sample environment map with the light direction
+            var sample_color = sample_environment(light_dir, source_mip).rgb;
+            sample_color = tonemap(sample_color);
+            
+            // Accumulate weighted sample, including geometric term
+            radiance += sample_color * NoL * G;
+            total_weight += NoL * G;
+        }
+    }
+    
+    // Normalize by total weight
+    if (total_weight > 0.0) {
+        radiance = radiance / total_weight;
+    }
+
+    // Reverse tonemap
+    radiance = reverse_tonemap(radiance);
+    
+    // Write result to output texture
+    textureStore(output_texture, coords, face, vec4f(radiance, 1.0));
+}
+
+// Calculate spherical coordinates using spiral pattern
+// and golden angle to get a uniform distribution
+fn uniform_sample_sphere(i: u32, normal: vec3f) -> vec3f {
+    // Get stratified sample index
+    let index = i % constants.sample_count;
+    
+    let golden_angle = 2.4;
+    let full_sphere = f32(constants.sample_count) * 2.0;
+    let z = 1.0 - (2.0 * f32(index) + 1.0) / full_sphere;
+    let r = sqrt(1.0 - z * z);
+    
+    let phi = f32(index) * golden_angle;
+    
+    // Create the direction vector
+    let dir_uniform = vec3f(
+        r * cos(phi),
+        r * sin(phi),
+        z
+    );
+
+    let tangent_frame = build_orthonormal_basis(normal);
+    return normalize(tangent_frame * dir_uniform);
+}
+
+@compute
+@workgroup_size(8, 8, 1)
+fn generate_irradiance_map(@builtin(global_invocation_id) global_id: vec3u) {
+    let size = textureDimensions(output_texture).xy;
+    let invSize = 1.0 / vec2f(size);
+    
+    let coords = vec2u(global_id.xy);
+    let face = global_id.z;
+    
+    if (any(coords >= size)) {
+        return;
+    }
+    
+    // Convert texture coordinates to direction vector
+    let uv = (vec2f(coords) + 0.5) * invSize;
+    let normal = sample_cube_dir(uv, face);
+    
+    var irradiance = vec3f(0.0);
+    var total_weight = 0.0;
+    
+    // Use uniform sampling on a hemisphere
+    for (var i = 0u; i < constants.sample_count; i++) {
+        // Get a uniform direction on unit sphere
+        var sample_dir = uniform_sample_sphere(i, normal);
+        
+        // Calculate the cosine weight (N·L)
+        let weight = max(dot(normal, sample_dir), 0.0);
+        
+        // Skip samples below horizon or at grazing angles
+        if (weight <= 0.001) {
+            continue;
+        }
+        
+        // Sample environment with level 0 (no mip)
+        var sample_color = sample_environment(sample_dir, 0.0).rgb;
+        
+        // Apply tonemapping to reduce fireflies
+        sample_color = tonemap(sample_color);
+        
+        // Accumulate the contribution
+        irradiance += sample_color * weight;
+        total_weight += weight;
+    }
+    
+    // Normalize by total weight
+    irradiance = irradiance / total_weight;
+    
+    // Reverse tonemap to restore HDR range
+    irradiance = reverse_tonemap(irradiance);
+    
+    // Write result to output texture
+    textureStore(output_texture, coords, face, vec4f(irradiance, 1.0));
+}

crates/bevy_pbr/src/light_probe/irradiance_volume.rs

@@ -127,7 +127,7 @@
 //!
 //! [Blender]: http://blender.org/
 //!
-//! [baking tool]: https://docs.blender.org/manual/en/latest/render/eevee/render_settings/indirect_lighting.html
+//! [baking tool]: https://docs.blender.org/manual/en/latest/render/eevee/light_probes/volume.html
 //!
 //! [`bevy-baked-gi`]: https://github.com/pcwalton/bevy-baked-gi
 //!

crates/bevy_pbr/src/light_probe/mod.rs

@@ -1,8 +1,11 @@
 //! Light probes for baked global illumination.
 
-use bevy_app::{App, Plugin};
-use bevy_asset::AssetId;
-use bevy_core_pipeline::core_3d::Camera3d;
+use bevy_app::{App, Plugin, Update};
+use bevy_asset::{embedded_asset, load_internal_binary_asset, AssetId, RenderAssetUsages};
+use bevy_core_pipeline::core_3d::{
+    graph::{Core3d, Node3d},
+    Camera3d,
+};
 use bevy_derive::{Deref, DerefMut};
 use bevy_ecs::{
     component::Component,
@@ -12,8 +15,8 @@ use bevy_ecs::{
     schedule::IntoScheduleConfigs,
     system::{Commands, Local, Query, Res, ResMut},
 };
-use bevy_image::Image;
-use bevy_light::{EnvironmentMapLight, LightProbe};
+use bevy_image::{CompressedImageFormats, Image, ImageSampler, ImageType};
+use bevy_light::{EnvironmentMapLight, GeneratedEnvironmentMapLight, LightProbe};
 use bevy_math::{Affine3A, FloatOrd, Mat4, Vec3A, Vec4};
 use bevy_platform::collections::HashMap;
 use bevy_render::{
@@ -21,24 +24,34 @@ use bevy_render::{
     load_shader_library,
     primitives::{Aabb, Frustum},
     render_asset::RenderAssets,
+    render_graph::RenderGraphExt,
     render_resource::{DynamicUniformBuffer, Sampler, ShaderType, TextureView},
     renderer::{RenderAdapter, RenderDevice, RenderQueue},
     settings::WgpuFeatures,
+    sync_component::SyncComponentPlugin,
     sync_world::RenderEntity,
     texture::{FallbackImage, GpuImage},
     view::ExtractedView,
-    Extract, ExtractSchedule, Render, RenderApp, RenderSystems,
+    Extract, ExtractSchedule, Render, RenderApp, RenderStartup, RenderSystems,
 };
 use bevy_transform::{components::Transform, prelude::GlobalTransform};
 use tracing::error;
 
 use core::{hash::Hash, ops::Deref};
 
-use crate::light_probe::environment_map::EnvironmentMapIds;
+use crate::{
+    generate::{
+        extract_generator_entities, generate_environment_map_light, init_generator_resources,
+        prepare_generator_bind_groups, prepare_intermediate_textures, GeneratorNode,
+        IrradianceMapNode, RadianceMapNode, SpdNode, STBN,
+    },
+    light_probe::environment_map::EnvironmentMapIds,
+};
 
 use self::irradiance_volume::IrradianceVolume;
 
 pub mod environment_map;
+pub mod generate;
 pub mod irradiance_volume;
 
 /// The maximum number of each type of light probe that each view will consider.
@@ -288,7 +301,25 @@ impl Plugin for LightProbePlugin {
         load_shader_library!(app, "environment_map.wgsl");
         load_shader_library!(app, "irradiance_volume.wgsl");
 
-        app.add_plugins(ExtractInstancesPlugin::<EnvironmentMapIds>::new());
+        embedded_asset!(app, "environment_filter.wgsl");
+        embedded_asset!(app, "spd.wgsl");
+        embedded_asset!(app, "copy_mip0.wgsl");
+
+        load_internal_binary_asset!(app, STBN, "stbn_vec2.png", |bytes, _: String| {
+            Image::from_buffer(
+                bytes,
+                ImageType::Extension("png"),
+                CompressedImageFormats::NONE,
+                false,
+                ImageSampler::Default,
+                RenderAssetUsages::RENDER_WORLD,
+            )
+            .expect("Failed to load spatio-temporal blue noise texture")
+        });
+
+        app.add_plugins(ExtractInstancesPlugin::<EnvironmentMapIds>::new())
+            .add_plugins(SyncComponentPlugin::<GeneratedEnvironmentMapLight>::default())
+            .add_systems(Update, generate_environment_map_light);
 
         let Some(render_app) = app.get_sub_app_mut(RenderApp) else {
             return;
@@ -297,14 +328,40 @@ impl Plugin for LightProbePlugin {
         render_app
             .init_resource::<LightProbesBuffer>()
             .init_resource::<EnvironmentMapUniformBuffer>()
+            .add_render_graph_node::<SpdNode>(Core3d, GeneratorNode::Mipmap)
+            .add_render_graph_node::<RadianceMapNode>(Core3d, GeneratorNode::Radiance)
+            .add_render_graph_node::<IrradianceMapNode>(Core3d, GeneratorNode::Irradiance)
+            .add_render_graph_edges(
+                Core3d,
+                (
+                    Node3d::EndPrepasses,
+                    GeneratorNode::Mipmap,
+                    GeneratorNode::Radiance,
+                    GeneratorNode::Irradiance,
+                    Node3d::StartMainPass,
+                ),
+            )
             .add_systems(ExtractSchedule, gather_environment_map_uniform)
             .add_systems(ExtractSchedule, gather_light_probes::<EnvironmentMapLight>)
             .add_systems(ExtractSchedule, gather_light_probes::<IrradianceVolume>)
+            .add_systems(
+                ExtractSchedule,
+                extract_generator_entities.after(generate_environment_map_light),
+            )
             .add_systems(
                 Render,
-                (upload_light_probes, prepare_environment_uniform_buffer)
+                prepare_generator_bind_groups.in_set(RenderSystems::PrepareBindGroups),
+            )
+            .add_systems(
+                Render,
+                (
+                    upload_light_probes,
+                    prepare_environment_uniform_buffer,
+                    prepare_intermediate_textures,
+                )
                     .in_set(RenderSystems::PrepareResources),
-            );
+            )
+            .add_systems(RenderStartup, init_generator_resources);
     }
 }
 

crates/bevy_pbr/src/light_probe/spd.wgsl

@@ -0,0 +1,398 @@
+// Ported from https://github.com/GPUOpen-LibrariesAndSDKs/FidelityFX-SDK/blob/c16b1d286b5b438b75da159ab51ff426bacea3d1/sdk/include/FidelityFX/gpu/spd/ffx_spd.h
+
+@group(0) @binding(0) var mip_0: texture_2d_array<f32>;
+@group(0) @binding(1) var mip_1: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(2) var mip_2: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(3) var mip_3: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(4) var mip_4: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(5) var mip_5: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(6) var mip_6: texture_storage_2d_array<rgba16float, read_write>;
+@group(0) @binding(7) var mip_7: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(8) var mip_8: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(9) var mip_9: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(10) var mip_10: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(11) var mip_11: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(12) var mip_12: texture_storage_2d_array<rgba16float, write>;
+@group(0) @binding(13) var sampler_linear_clamp: sampler;
+@group(0) @binding(14) var<uniform> constants: Constants;
+struct Constants { mips: u32, inverse_input_size: vec2f }
+
+var<workgroup> spd_intermediate_r: array<array<f32, 16>, 16>;
+var<workgroup> spd_intermediate_g: array<array<f32, 16>, 16>;
+var<workgroup> spd_intermediate_b: array<array<f32, 16>, 16>;
+var<workgroup> spd_intermediate_a: array<array<f32, 16>, 16>;
+
+@compute
+@workgroup_size(256, 1, 1)
+fn spd_downsample_first(
+    @builtin(workgroup_id) workgroup_id: vec3u,
+    @builtin(local_invocation_index) local_invocation_index: u32,
+#ifdef SUBGROUP_SUPPORT
+    @builtin(subgroup_invocation_id) subgroup_invocation_id: u32,
+#endif
+) {
+#ifndef SUBGROUP_SUPPORT
+    let subgroup_invocation_id = 0u;
+#endif
+
+    let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
+    let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
+    let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
+
+    spd_downsample_mips_0_1(x, y, workgroup_id.xy, local_invocation_index, constants.mips, workgroup_id.z, subgroup_invocation_id);
+
+    spd_downsample_next_four(x, y, workgroup_id.xy, local_invocation_index, 2u, constants.mips, workgroup_id.z, subgroup_invocation_id);
+}
+
+// TODO: Once wgpu supports globallycoherent buffers, make it actually a single pass
+@compute
+@workgroup_size(256, 1, 1)
+fn spd_downsample_second(
+    @builtin(workgroup_id) workgroup_id: vec3u,
+    @builtin(local_invocation_index) local_invocation_index: u32,
+#ifdef SUBGROUP_SUPPORT
+    @builtin(subgroup_invocation_id) subgroup_invocation_id: u32,
+#endif
+) {
+#ifndef SUBGROUP_SUPPORT
+    let subgroup_invocation_id = 0u;
+#endif
+
+    let sub_xy = remap_for_wave_reduction(local_invocation_index % 64u);
+    let x = sub_xy.x + 8u * ((local_invocation_index >> 6u) % 2u);
+    let y = sub_xy.y + 8u * (local_invocation_index >> 7u);
+
+    spd_downsample_mips_6_7(x, y, constants.mips, workgroup_id.z);
+
+    spd_downsample_next_four(x, y, vec2(0u), local_invocation_index, 8u, constants.mips, workgroup_id.z, subgroup_invocation_id);
+}
+
+fn spd_downsample_mips_0_1(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, mips: u32, slice: u32, subgroup_invocation_id: u32) {
+    var v: array<vec4f, 4>;
+
+    var tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u);
+    var pix = (workgroup_id * 32u) + vec2(x, y);
+    v[0] = spd_reduce_load_source_image(tex, slice);
+    spd_store(pix, v[0], 0u, slice);
+
+    tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u);
+    pix = (workgroup_id * 32u) + vec2(x + 16u, y);
+    v[1] = spd_reduce_load_source_image(tex, slice);
+    spd_store(pix, v[1], 0u, slice);
+
+    tex = (workgroup_id * 64u) + vec2(x * 2u, y * 2u + 32u);
+    pix = (workgroup_id * 32u) + vec2(x, y + 16u);
+    v[2] = spd_reduce_load_source_image(tex, slice);
+    spd_store(pix, v[2], 0u, slice);
+
+    tex = (workgroup_id * 64u) + vec2(x * 2u + 32u, y * 2u + 32u);
+    pix = (workgroup_id * 32u) + vec2(x + 16u, y + 16u);
+    v[3] = spd_reduce_load_source_image(tex, slice);
+    spd_store(pix, v[3], 0u, slice);
+
+    if mips <= 1u { return; }
+
+#ifdef SUBGROUP_SUPPORT
+    v[0] = spd_reduce_quad(v[0], subgroup_invocation_id);
+    v[1] = spd_reduce_quad(v[1], subgroup_invocation_id);
+    v[2] = spd_reduce_quad(v[2], subgroup_invocation_id);
+    v[3] = spd_reduce_quad(v[3], subgroup_invocation_id);
+
+    if local_invocation_index % 4u == 0u {
+        spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u), v[0], 1u, slice);
+        spd_store_intermediate(x / 2u, y / 2u, v[0]);
+
+        spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u), v[1], 1u, slice);
+        spd_store_intermediate(x / 2u + 8u, y / 2u, v[1]);
+
+        spd_store((workgroup_id * 16u) + vec2(x / 2u, y / 2u + 8u), v[2], 1u, slice);
+        spd_store_intermediate(x / 2u, y / 2u + 8u, v[2]);
+
+        spd_store((workgroup_id * 16u) + vec2(x / 2u + 8u, y / 2u + 8u), v[3], 1u, slice);
+        spd_store_intermediate(x / 2u + 8u, y / 2u + 8u, v[3]);
+    }
+#else
+    for (var i = 0u; i < 4u; i++) {
+        spd_store_intermediate(x, y, v[i]);
+        workgroupBarrier();
+        if local_invocation_index < 64u {
+            v[i] = spd_reduce_intermediate(
+                vec2(x * 2u + 0u, y * 2u + 0u),
+                vec2(x * 2u + 1u, y * 2u + 0u),
+                vec2(x * 2u + 0u, y * 2u + 1u),
+                vec2(x * 2u + 1u, y * 2u + 1u),
+            );
+            spd_store(vec2(workgroup_id * 16) + vec2(x + (i % 2u) * 8u, y + (i / 2u) * 8u), v[i], 1u, slice);
+        }
+        workgroupBarrier();
+    }
+
+    if local_invocation_index < 64u {
+        spd_store_intermediate(x + 0u, y + 0u, v[0]);
+        spd_store_intermediate(x + 8u, y + 0u, v[1]);
+        spd_store_intermediate(x + 0u, y + 8u, v[2]);
+        spd_store_intermediate(x + 8u, y + 8u, v[3]);
+    }
+#endif
+}
+
+fn spd_downsample_next_four(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, mips: u32, slice: u32, subgroup_invocation_id: u32) {
+    if mips <= base_mip { return; }
+    workgroupBarrier();
+    spd_downsample_mip_2(x, y, workgroup_id, local_invocation_index, base_mip, slice, subgroup_invocation_id);
+
+    if mips <= base_mip + 1u { return; }
+    workgroupBarrier();
+    spd_downsample_mip_3(x, y, workgroup_id, local_invocation_index, base_mip + 1u, slice, subgroup_invocation_id);
+
+    if mips <= base_mip + 2u { return; }
+    workgroupBarrier();
+    spd_downsample_mip_4(x, y, workgroup_id, local_invocation_index, base_mip + 2u, slice, subgroup_invocation_id);
+
+    if mips <= base_mip + 3u { return; }
+    workgroupBarrier();
+    spd_downsample_mip_5(x, y, workgroup_id, local_invocation_index, base_mip + 3u, slice, subgroup_invocation_id);
+}
+
+fn spd_downsample_mip_2(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+#ifdef SUBGROUP_SUPPORT
+    var v = spd_load_intermediate(x, y);
+    v = spd_reduce_quad(v, subgroup_invocation_id);
+    if local_invocation_index % 4u == 0u {
+        spd_store((workgroup_id * 8u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
+        spd_store_intermediate(x + (y / 2u) % 2u, y, v);
+    }
+#else
+    if local_invocation_index < 64u {
+        let v = spd_reduce_intermediate(
+            vec2(x * 2u + 0u, y * 2u + 0u),
+            vec2(x * 2u + 1u, y * 2u + 0u),
+            vec2(x * 2u + 0u, y * 2u + 1u),
+            vec2(x * 2u + 1u, y * 2u + 1u),
+        );
+        spd_store((workgroup_id * 8u) + vec2(x, y), v, base_mip, slice);
+        spd_store_intermediate(x * 2u + y % 2u, y * 2u, v);
+    }
+#endif
+}
+
+fn spd_downsample_mip_3(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+#ifdef SUBGROUP_SUPPORT
+    if local_invocation_index < 64u {
+        var v = spd_load_intermediate(x * 2u + y % 2u, y * 2u);
+        v = spd_reduce_quad(v, subgroup_invocation_id);
+        if local_invocation_index % 4u == 0u {
+            spd_store((workgroup_id * 4u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
+            spd_store_intermediate(x * 2u + y / 2u, y * 2u, v);
+        }
+    }
+#else
+    if local_invocation_index < 16u {
+        let v = spd_reduce_intermediate(
+            vec2(x * 4u + 0u + 0u, y * 4u + 0u),
+            vec2(x * 4u + 2u + 0u, y * 4u + 0u),
+            vec2(x * 4u + 0u + 1u, y * 4u + 2u),
+            vec2(x * 4u + 2u + 1u, y * 4u + 2u),
+        );
+        spd_store((workgroup_id * 4u) + vec2(x, y), v, base_mip, slice);
+        spd_store_intermediate(x * 4u + y, y * 4u, v);
+    }
+#endif
+}
+
+fn spd_downsample_mip_4(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+#ifdef SUBGROUP_SUPPORT
+    if local_invocation_index < 16u {
+        var v = spd_load_intermediate(x * 4u + y, y * 4u);
+        v = spd_reduce_quad(v, subgroup_invocation_id);
+        if local_invocation_index % 4u == 0u {
+            spd_store((workgroup_id * 2u) + vec2(x / 2u, y / 2u), v, base_mip, slice);
+            spd_store_intermediate(x / 2u + y, 0u, v);
+        }
+    }
+#else
+    if local_invocation_index < 4u {
+        let v = spd_reduce_intermediate(
+            vec2(x * 8u + 0u + 0u + y * 2u, y * 8u + 0u),
+            vec2(x * 8u + 4u + 0u + y * 2u, y * 8u + 0u),
+            vec2(x * 8u + 0u + 1u + y * 2u, y * 8u + 4u),
+            vec2(x * 8u + 4u + 1u + y * 2u, y * 8u + 4u),
+        );
+        spd_store((workgroup_id * 2u) + vec2(x, y), v, base_mip, slice);
+        spd_store_intermediate(x + y * 2u, 0u, v);
+    }
+#endif
+}
+
+fn spd_downsample_mip_5(x: u32, y: u32, workgroup_id: vec2u, local_invocation_index: u32, base_mip: u32, slice: u32, subgroup_invocation_id: u32) {
+#ifdef SUBGROUP_SUPPORT
+    if local_invocation_index < 4u {
+        var v = spd_load_intermediate(local_invocation_index, 0u);
+        v = spd_reduce_quad(v, subgroup_invocation_id);
+        if local_invocation_index % 4u == 0u {
+            spd_store(workgroup_id, v, base_mip, slice);
+        }
+    }
+#else
+    if local_invocation_index < 1u {
+        let v = spd_reduce_intermediate(vec2(0u, 0u), vec2(1u, 0u), vec2(2u, 0u), vec2(3u, 0u));
+        spd_store(workgroup_id, v, base_mip, slice);
+    }
+#endif
+}
+
+fn spd_downsample_mips_6_7(x: u32, y: u32, mips: u32, slice: u32) {
+    var tex = vec2(x * 4u + 0u, y * 4u + 0u);
+    var pix = vec2(x * 2u + 0u, y * 2u + 0u);
+    let v0 = spd_reduce_load_4(
+        vec2(x * 4u + 0u, y * 4u + 0u),
+        vec2(x * 4u + 1u, y * 4u + 0u),
+        vec2(x * 4u + 0u, y * 4u + 1u),
+        vec2(x * 4u + 1u, y * 4u + 1u),
+        slice
+    );
+    spd_store(pix, v0, 6u, slice);
+
+    tex = vec2(x * 4u + 2u, y * 4u + 0u);
+    pix = vec2(x * 2u + 1u, y * 2u + 0u);
+    let v1 = spd_reduce_load_4(
+        vec2(x * 4u + 2u, y * 4u + 0u),
+        vec2(x * 4u + 3u, y * 4u + 0u),
+        vec2(x * 4u + 2u, y * 4u + 1u),
+        vec2(x * 4u + 3u, y * 4u + 1u),
+        slice
+    );
+    spd_store(pix, v1, 6u, slice);
+
+    tex = vec2(x * 4u + 0u, y * 4u + 2u);
+    pix = vec2(x * 2u + 0u, y * 2u + 1u);
+    let v2 = spd_reduce_load_4(
+        vec2(x * 4u + 0u, y * 4u + 2u),
+        vec2(x * 4u + 1u, y * 4u + 2u),
+        vec2(x * 4u + 0u, y * 4u + 3u),
+        vec2(x * 4u + 1u, y * 4u + 3u),
+        slice
+    );
+    spd_store(pix, v2, 6u, slice);
+
+    tex = vec2(x * 4u + 2u, y * 4u + 2u);
+    pix = vec2(x * 2u + 1u, y * 2u + 1u);
+    let v3 = spd_reduce_load_4(
+        vec2(x * 4u + 2u, y * 4u + 2u),
+        vec2(x * 4u + 3u, y * 4u + 2u),
+        vec2(x * 4u + 2u, y * 4u + 3u),
+        vec2(x * 4u + 3u, y * 4u + 3u),
+        slice
+    );
+    spd_store(pix, v3, 6u, slice);
+
+    if mips < 7u { return; }
+
+    let v = spd_reduce_4(v0, v1, v2, v3);
+    spd_store(vec2(x, y), v, 7u, slice);
+    spd_store_intermediate(x, y, v);
+}
+
+fn remap_for_wave_reduction(a: u32) -> vec2u {
+    // This function maps linear thread IDs to 2D coordinates in a special pattern
+    // to ensure that neighboring threads process neighboring pixels
+    // For example, this transforms linear thread IDs 0,1,2,3 into a 2×2 square
+    
+    // Extract bits to form the X and Y coordinates
+    let x = insertBits(extractBits(a, 2u, 3u), a, 0u, 1u);
+    let y = insertBits(extractBits(a, 3u, 3u), extractBits(a, 1u, 2u), 0u, 2u);
+    
+    return vec2u(x, y);
+}
+
+fn spd_reduce_load_source_image(uv: vec2u, slice: u32) -> vec4f {
+    let texture_coord = (vec2f(uv) + 0.5) * constants.inverse_input_size;
+    
+    let result = textureSampleLevel(mip_0, sampler_linear_clamp, texture_coord, slice, 0.0);
+
+#ifdef SRGB_CONVERSION
+    return vec4(
+        srgb_from_linear(result.r),
+        srgb_from_linear(result.g),
+        srgb_from_linear(result.b),
+        result.a
+    );
+#else
+    return result;
+#endif
+
+}
+
+fn spd_store(pix: vec2u, value: vec4f, mip: u32, slice: u32) {
+    if mip >= constants.mips { return; }
+    switch mip {
+        case 0u: { textureStore(mip_1, pix, slice, value); }
+        case 1u: { textureStore(mip_2, pix, slice, value); }
+        case 2u: { textureStore(mip_3, pix, slice, value); }
+        case 3u: { textureStore(mip_4, pix, slice, value); }
+        case 4u: { textureStore(mip_5, pix, slice, value); }
+        case 5u: { textureStore(mip_6, pix, slice, value); }
+        case 6u: { textureStore(mip_7, pix, slice, value); }
+        case 7u: { textureStore(mip_8, pix, slice, value); }
+        case 8u: { textureStore(mip_9, pix, slice, value); }
+        case 9u: { textureStore(mip_10, pix, slice, value); }
+        case 10u: { textureStore(mip_11, pix, slice, value); }
+        case 11u: { textureStore(mip_12, pix, slice, value); }
+        default: {}
+    }
+}
+
+fn spd_store_intermediate(x: u32, y: u32, value: vec4f) {
+    spd_intermediate_r[x][y] = value.x;
+    spd_intermediate_g[x][y] = value.y;
+    spd_intermediate_b[x][y] = value.z;
+    spd_intermediate_a[x][y] = value.w;
+}
+
+fn spd_load_intermediate(x: u32, y: u32) -> vec4f {
+    return vec4(spd_intermediate_r[x][y], spd_intermediate_g[x][y], spd_intermediate_b[x][y], spd_intermediate_a[x][y]);
+}
+
+fn spd_reduce_intermediate(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u) -> vec4f {
+    let v0 = spd_load_intermediate(i0.x, i0.y);
+    let v1 = spd_load_intermediate(i1.x, i1.y);
+    let v2 = spd_load_intermediate(i2.x, i2.y);
+    let v3 = spd_load_intermediate(i3.x, i3.y);
+    return spd_reduce_4(v0, v1, v2, v3);
+}
+
+fn spd_reduce_load_4(i0: vec2u, i1: vec2u, i2: vec2u, i3: vec2u, slice: u32) -> vec4f {
+    let v0 = textureLoad(mip_6, i0, slice);
+    let v1 = textureLoad(mip_6, i1, slice);
+    let v2 = textureLoad(mip_6, i2, slice);
+    let v3 = textureLoad(mip_6, i3, slice);
+    return spd_reduce_4(v0, v1, v2, v3);
+}
+
+fn spd_reduce_4(v0: vec4f, v1: vec4f, v2: vec4f, v3: vec4f) -> vec4f {
+    return (v0 + v1 + v2 + v3) * 0.25;
+}
+
+#ifdef SUBGROUP_SUPPORT
+fn spd_reduce_quad(v: vec4f, subgroup_invocation_id: u32) -> vec4f {
+    let quad = subgroup_invocation_id & (~0x3u);
+    let v0 = v;
+    let v1 = subgroupBroadcast(v, quad | 1u);
+    let v2 = subgroupBroadcast(v, quad | 2u);
+    let v3 = subgroupBroadcast(v, quad | 3u);
+    return spd_reduce_4(v0, v1, v2, v3);
+
+    // TODO: Use subgroup quad operations once wgpu supports them
+    // let v0 = v;
+    // let v1 = quadSwapX(v);
+    // let v2 = quadSwapY(v);
+    // let v3 = quadSwapDiagonal(v);
+    // return spd_reduce_4(v0, v1, v2, v3);
+}
+#endif
+
+fn srgb_from_linear(value: f32) -> f32 {
+    let j = vec3(0.0031308 * 12.92, 12.92, 1.0 / 2.4);
+    let k = vec2(1.055, -0.055);
+    return clamp(j.x, value * j.y, pow(value, j.z) * k.x + k.y);
+}

crates/bevy_pbr/src/render/utils.wgsl

@@ -68,6 +68,16 @@ fn octahedral_decode_signed(v: vec2<f32>) -> vec3<f32> {
     return normalize(n);
 }
 
+// https://jcgt.org/published/0006/01/01/paper.pdf
+fn build_orthonormal_basis(normal: vec3<f32>) -> mat3x3<f32> {
+    let sign = select(-1.0, 1.0, normal.z >= 0.0);
+    let a = -1.0 / (sign + normal.z);
+    let b = normal.x * normal.y * a;
+    let tangent = vec3(1.0 + sign * normal.x * normal.x * a, sign * b, -sign * normal.x);
+    let bitangent = vec3(b, sign + normal.y * normal.y * a, -normal.y);
+    return mat3x3(tangent, bitangent, normal);
+}
+
 // https://blog.demofox.org/2022/01/01/interleaved-gradient-noise-a-different-kind-of-low-discrepancy-sequence
 fn interleaved_gradient_noise(pixel_coordinates: vec2<f32>, frame: u32) -> f32 {
     let xy = pixel_coordinates + 5.588238 * f32(frame % 64u);

crates/bevy_solari/src/scene/sampling.wgsl

@@ -1,6 +1,6 @@
 #define_import_path bevy_solari::sampling
 
-#import bevy_pbr::utils::{rand_f, rand_vec2f, rand_range_u}
+#import bevy_pbr::utils::{rand_f, rand_vec2f, rand_range_u, build_orthonormal_basis}
 #import bevy_render::maths::{PI, PI_2}
 #import bevy_solari::scene_bindings::{trace_ray, RAY_T_MIN, RAY_T_MAX, light_sources, directional_lights, LIGHT_SOURCE_KIND_DIRECTIONAL, resolve_triangle_data_full}
 
@@ -208,13 +208,3 @@ fn triangle_barycentrics(random: vec2<f32>) -> vec3<f32> {
     if barycentrics.x + barycentrics.y > 1.0 { barycentrics = 1.0 - barycentrics; }
     return vec3(1.0 - barycentrics.x - barycentrics.y, barycentrics);
 }
-
-// https://jcgt.org/published/0006/01/01/paper.pdf
-fn build_orthonormal_basis(normal: vec3<f32>) -> mat3x3<f32> {
-    let sign = select(-1.0, 1.0, normal.z >= 0.0);
-    let a = -1.0 / (sign + normal.z);
-    let b = normal.x * normal.y * a;
-    let tangent = vec3(1.0 + sign * normal.x * normal.x * a, sign * b, -sign * normal.x);
-    let bitangent = vec3(b, sign + normal.y * normal.y * a, -normal.y);
-    return mat3x3(tangent, bitangent, normal);
-}

examples/3d/reflection_probes.rs

@@ -1,12 +1,17 @@
 //! This example shows how to place reflection probes in the scene.
 //!
-//! Press Space to switch between no reflections, environment map reflections
-//! (i.e. the skybox only, not the cubes), and a full reflection probe that
-//! reflects the skybox and the cubes. Press Enter to pause rotation.
+//! Press Space to cycle through the reflection modes: an environment-map mode
+//! that shows only the skybox, a reflection-probe mode that reflects both the
+//! skybox and the cubes, and a generated environment map mode that filters an
+//! unfiltered cubemap on the GPU. Press Enter to pause or resume rotation.
 //!
 //! Reflection probes don't work on WebGL 2 or WebGPU.
 
-use bevy::{core_pipeline::Skybox, prelude::*, render::view::Hdr};
+use bevy::{
+    core_pipeline::{tonemapping::Tonemapping, Skybox},
+    prelude::*,
+    render::{render_resource::TextureUsages, view::Hdr},
+};
 
 use std::{
     f32::consts::PI,
@@ -25,37 +30,33 @@ struct AppStatus {
     reflection_mode: ReflectionMode,
     // Whether the user has requested the scene to rotate.
     rotating: bool,
+    // The current roughness of the central sphere
+    sphere_roughness: f32,
 }
 
 // Which environment maps the user has requested to display.
-#[derive(Clone, Copy)]
+#[derive(Clone, Copy, PartialEq)]
 enum ReflectionMode {
-    // No environment maps are shown.
-    None = 0,
     // Only a world environment map is shown.
-    EnvironmentMap = 1,
+    EnvironmentMap = 0,
     // Both a world environment map and a reflection probe are present. The
     // reflection probe is shown in the sphere.
-    ReflectionProbe = 2,
+    ReflectionProbe = 1,
+    // A generated environment map is shown.
+    GeneratedEnvironmentMap = 2,
 }
 
 // The various reflection maps.
 #[derive(Resource)]
 struct Cubemaps {
-    // The blurry diffuse cubemap. This is used for both the world environment
-    // map and the reflection probe. (In reality you wouldn't do this, but this
-    // reduces complexity of this example a bit.)
-    diffuse: Handle<Image>,
+    // The blurry diffuse cubemap that reflects the world, but not the cubes.
+    diffuse_environment_map: Handle<Image>,
 
-    // The specular cubemap that reflects the world, but not the cubes.
+    // The specular cubemap mip chain that reflects the world, but not the cubes.
     specular_environment_map: Handle<Image>,
 
-    // The specular cubemap that reflects both the world and the cubes.
+    // The specular cubemap mip chain that reflects both the world and the cubes.
     specular_reflection_probe: Handle<Image>,
-
-    // The skybox cubemap image. This is almost the same as
-    // `specular_environment_map`.
-    skybox: Handle<Image>,
 }
 
 fn main() {
@@ -68,6 +69,7 @@ fn main() {
         .add_systems(PreUpdate, add_environment_map_to_camera)
         .add_systems(Update, change_reflection_type)
         .add_systems(Update, toggle_rotation)
+        .add_systems(Update, change_sphere_roughness)
         .add_systems(
             Update,
             rotate_camera
@@ -75,6 +77,7 @@ fn main() {
                 .after(change_reflection_type),
         )
         .add_systems(Update, update_text.after(rotate_camera))
+        .add_systems(Update, setup_environment_map_usage)
         .run();
 }
 
@@ -89,7 +92,7 @@ fn setup(
 ) {
     spawn_scene(&mut commands, &asset_server);
     spawn_camera(&mut commands);
-    spawn_sphere(&mut commands, &mut meshes, &mut materials);
+    spawn_sphere(&mut commands, &mut meshes, &mut materials, &app_status);
     spawn_reflection_probe(&mut commands, &cubemaps);
     spawn_text(&mut commands, &app_status);
 }
@@ -105,8 +108,9 @@ fn spawn_scene(commands: &mut Commands, asset_server: &AssetServer) {
 fn spawn_camera(commands: &mut Commands) {
     commands.spawn((
         Camera3d::default(),
-        Transform::from_xyz(-6.483, 0.325, 4.381).looking_at(Vec3::ZERO, Vec3::Y),
         Hdr,
+        Tonemapping::AcesFitted,
+        Transform::from_xyz(-6.483, 0.325, 4.381).looking_at(Vec3::ZERO, Vec3::Y),
     ));
 }
 
@@ -115,6 +119,7 @@ fn spawn_sphere(
     commands: &mut Commands,
     meshes: &mut Assets<Mesh>,
     materials: &mut Assets<StandardMaterial>,
+    app_status: &AppStatus,
 ) {
     // Create a sphere mesh.
     let sphere_mesh = meshes.add(Sphere::new(1.0).mesh().ico(7).unwrap());
@@ -123,11 +128,12 @@ fn spawn_sphere(
     commands.spawn((
         Mesh3d(sphere_mesh.clone()),
         MeshMaterial3d(materials.add(StandardMaterial {
-            base_color: Srgba::hex("#ffd891").unwrap().into(),
+            base_color: Srgba::hex("#ffffff").unwrap().into(),
             metallic: 1.0,
-            perceptual_roughness: 0.0,
+            perceptual_roughness: app_status.sphere_roughness,
             ..StandardMaterial::default()
         })),
+        SphereMaterial,
     ));
 }
 
@@ -136,7 +142,7 @@ fn spawn_reflection_probe(commands: &mut Commands, cubemaps: &Cubemaps) {
     commands.spawn((
         LightProbe,
         EnvironmentMapLight {
-            diffuse_map: cubemaps.diffuse.clone(),
+            diffuse_map: cubemaps.diffuse_environment_map.clone(),
             specular_map: cubemaps.specular_reflection_probe.clone(),
             intensity: 5000.0,
             ..default()
@@ -173,7 +179,7 @@ fn add_environment_map_to_camera(
             .entity(camera_entity)
             .insert(create_camera_environment_map_light(&cubemaps))
             .insert(Skybox {
-                image: cubemaps.skybox.clone(),
+                image: cubemaps.specular_environment_map.clone(),
                 brightness: 5000.0,
                 ..default()
             });
@@ -194,30 +200,45 @@ fn change_reflection_type(
         return;
     }
 
-    // Switch reflection mode.
+    // Advance to the next reflection mode.
     app_status.reflection_mode =
         ReflectionMode::try_from((app_status.reflection_mode as u32 + 1) % 3).unwrap();
 
-    // Add or remove the light probe.
+    // Remove light probes
     for light_probe in light_probe_query.iter() {
         commands.entity(light_probe).despawn();
     }
     match app_status.reflection_mode {
-        ReflectionMode::None | ReflectionMode::EnvironmentMap => {}
+        ReflectionMode::EnvironmentMap => {}
         ReflectionMode::ReflectionProbe => spawn_reflection_probe(&mut commands, &cubemaps),
+        ReflectionMode::GeneratedEnvironmentMap => {}
     }
 
-    // Add or remove the environment map from the camera.
+    // Update the environment-map components on the camera entity/entities
     for camera in camera_query.iter() {
+        // Remove any existing environment-map components
+        commands
+            .entity(camera)
+            .remove::<(EnvironmentMapLight, GeneratedEnvironmentMapLight)>();
+
         match app_status.reflection_mode {
-            ReflectionMode::None => {
-                commands.entity(camera).remove::<EnvironmentMapLight>();
-            }
+            // A baked or reflection-probe environment map
             ReflectionMode::EnvironmentMap | ReflectionMode::ReflectionProbe => {
                 commands
                     .entity(camera)
                     .insert(create_camera_environment_map_light(&cubemaps));
             }
+
+            // GPU-filtered environment map generated at runtime
+            ReflectionMode::GeneratedEnvironmentMap => {
+                commands
+                    .entity(camera)
+                    .insert(GeneratedEnvironmentMapLight {
+                        environment_map: cubemaps.specular_environment_map.clone(),
+                        intensity: 5000.0,
+                        ..default()
+                    });
+            }
         }
     }
 }
@@ -241,9 +262,9 @@ impl TryFrom<u32> for ReflectionMode {
 
     fn try_from(value: u32) -> Result<Self, Self::Error> {
         match value {
-            0 => Ok(ReflectionMode::None),
-            1 => Ok(ReflectionMode::EnvironmentMap),
-            2 => Ok(ReflectionMode::ReflectionProbe),
+            0 => Ok(ReflectionMode::EnvironmentMap),
+            1 => Ok(ReflectionMode::ReflectionProbe),
+            2 => Ok(ReflectionMode::GeneratedEnvironmentMap),
             _ => Err(()),
         }
     }
@@ -252,9 +273,9 @@ impl TryFrom<u32> for ReflectionMode {
 impl Display for ReflectionMode {
     fn fmt(&self, formatter: &mut Formatter<'_>) -> FmtResult {
         let text = match *self {
-            ReflectionMode::None => "No reflections",
             ReflectionMode::EnvironmentMap => "Environment map",
             ReflectionMode::ReflectionProbe => "Reflection probe",
+            ReflectionMode::GeneratedEnvironmentMap => "Generated environment map",
         };
         formatter.write_str(text)
     }
@@ -271,8 +292,11 @@ impl AppStatus {
         };
 
         format!(
-            "{}\n{}\n{}",
-            self.reflection_mode, rotation_help_text, REFLECTION_MODE_HELP_TEXT
+            "{}\n{}\nRoughness: {:.2}\n{}\nUp/Down arrows to change roughness",
+            self.reflection_mode,
+            rotation_help_text,
+            self.sphere_roughness,
+            REFLECTION_MODE_HELP_TEXT
         )
         .into()
     }
@@ -282,7 +306,7 @@ impl AppStatus {
 // probe is applicable to a mesh.
 fn create_camera_environment_map_light(cubemaps: &Cubemaps) -> EnvironmentMapLight {
     EnvironmentMapLight {
-        diffuse_map: cubemaps.diffuse.clone(),
+        diffuse_map: cubemaps.diffuse_environment_map.clone(),
         specular_map: cubemaps.specular_environment_map.clone(),
         intensity: 5000.0,
         ..default()
@@ -311,17 +335,25 @@ fn rotate_camera(
 // Loads the cubemaps from the assets directory.
 impl FromWorld for Cubemaps {
     fn from_world(world: &mut World) -> Self {
-        // Just use the specular map for the skybox since it's not too blurry.
-        // In reality you wouldn't do this--you'd use a real skybox texture--but
-        // reusing the textures like this saves space in the Bevy repository.
-        let specular_map = world.load_asset("environment_maps/pisa_specular_rgb9e5_zstd.ktx2");
-
         Cubemaps {
-            diffuse: world.load_asset("environment_maps/pisa_diffuse_rgb9e5_zstd.ktx2"),
+            diffuse_environment_map: world
+                .load_asset("environment_maps/pisa_diffuse_rgb9e5_zstd.ktx2"),
+            specular_environment_map: world
+                .load_asset("environment_maps/pisa_specular_rgb9e5_zstd.ktx2"),
             specular_reflection_probe: world
                 .load_asset("environment_maps/cubes_reflection_probe_specular_rgb9e5_zstd.ktx2"),
-            specular_environment_map: specular_map.clone(),
-            skybox: specular_map,
+        }
+    }
+}
+
+fn setup_environment_map_usage(cubemaps: Res<Cubemaps>, mut images: ResMut<Assets<Image>>) {
+    if let Some(image) = images.get_mut(&cubemaps.specular_environment_map) {
+        if !image
+            .texture_descriptor
+            .usage
+            .contains(TextureUsages::COPY_SRC)
+        {
+            image.texture_descriptor.usage |= TextureUsages::COPY_SRC;
         }
     }
 }
@@ -331,6 +363,39 @@ impl Default for AppStatus {
         Self {
             reflection_mode: ReflectionMode::ReflectionProbe,
             rotating: true,
+            sphere_roughness: 0.0,
+        }
+    }
+}
+
+#[derive(Component)]
+struct SphereMaterial;
+
+// A system that changes the sphere's roughness with up/down arrow keys
+fn change_sphere_roughness(
+    keyboard: Res<ButtonInput<KeyCode>>,
+    mut app_status: ResMut<AppStatus>,
+    mut materials: ResMut<Assets<StandardMaterial>>,
+    sphere_query: Query<&MeshMaterial3d<StandardMaterial>, With<SphereMaterial>>,
+) {
+    let roughness_delta = if keyboard.pressed(KeyCode::ArrowUp) {
+        0.01 // Decrease roughness
+    } else if keyboard.pressed(KeyCode::ArrowDown) {
+        -0.01 // Increase roughness
+    } else {
+        0.0 // No change
+    };
+
+    if roughness_delta != 0.0 {
+        // Update the app status
+        app_status.sphere_roughness =
+            (app_status.sphere_roughness + roughness_delta).clamp(0.0, 1.0);
+
+        // Update the sphere material
+        for material_handle in sphere_query.iter() {
+            if let Some(material) = materials.get_mut(&material_handle.0) {
+                material.perceptual_roughness = app_status.sphere_roughness;
+            }
         }
     }
 }