diff --git a/blade-helpers/src/hud.rs b/blade-helpers/src/hud.rs index e8d0fea8..ff4cf5a5 100644 --- a/blade-helpers/src/hud.rs +++ b/blade-helpers/src/hud.rs @@ -22,7 +22,7 @@ impl ExposeHud for blade_render::RayConfig { egui::widgets::Slider::new(&mut self.spatial_tap_history, 0..=50) .text("Spatial tap history"), ); - ui.checkbox(&mut self.spatial_jitter, "Spatial jittering"); + ui.add(egui::widgets::Slider::new(&mut self.group_mixer, 1..=10).text("Group mixer")); ui.add( egui::widgets::Slider::new(&mut self.spatial_min_distance, 1..=10) .text("Spatial minimum distance (px)"), diff --git a/blade-render/Cargo.toml b/blade-render/Cargo.toml index 4d2bc38f..7017bfd9 100644 --- a/blade-render/Cargo.toml +++ b/blade-render/Cargo.toml @@ -39,7 +39,6 @@ glam = { workspace = true } log = { workspace = true } mikktspace = { package = "bevy_mikktspace", version = "0.12", optional = true } mint = { workspace = true } -nanorand = { workspace = true, features = ["wyrand"] } profiling = { workspace = true } slab = { workspace = true, optional = true } strum = { workspace = true } diff --git a/blade-render/code/color.inc.wgsl b/blade-render/code/color.inc.wgsl new file mode 100644 index 00000000..8e84cfda --- /dev/null +++ b/blade-render/code/color.inc.wgsl @@ -0,0 +1,19 @@ +fn hsv_to_rgb(h: f32, s: f32, v: f32) -> vec3 { + let c = v * s; + let x = c * (1.0 - abs((h / 60.0) % 2.0 - 1.0)); + var q = vec3(v - c); + if (h < 60.0) { + q.r += c; q.g += x; + } else if (h < 120.0) { + q.g += c; q.r += x; + } else if (h < 180.0) { + q.g += c; q.b += x; + } else if (h < 240.0) { + q.b += c; q.g += x; + } else if (h < 300.0) { + q.b += c; q.r += x; + } else { + q.r += c; q.b += x; + } + return q; +} diff --git a/blade-render/code/ray-trace.wgsl b/blade-render/code/ray-trace.wgsl index db06b5e4..416a1acb 100644 --- a/blade-render/code/ray-trace.wgsl +++ b/blade-render/code/ray-trace.wgsl @@ -1,3 +1,4 @@ +#include "color.inc.wgsl" #include "quaternion.inc.wgsl" #include "random.inc.wgsl" #include "env-importance.inc.wgsl" @@ -26,6 +27,7 @@ const DECOUPLED_SHADING: bool = false; //TODO: crashes on AMD 6850U if `GROUP_SIZE_TOTAL` > 32 const GROUP_SIZE: vec2 = vec2(8, 4); const GROUP_SIZE_TOTAL: u32 = GROUP_SIZE.x * GROUP_SIZE.y; +const GROUP_VISUALIZE: bool = false; struct MainParams { frame_index: u32, @@ -38,7 +40,6 @@ struct MainParams { spatial_min_distance: i32, t_start: f32, use_motion_vectors: u32, - grid_offset: vec2, grid_scale: vec2, } @@ -231,7 +232,6 @@ fn thread_index_to_coord(thread_index: u32, group_id: vec3) -> vec2 { let cluster_offset = group_id.xy - cluster_id * parameters.grid_scale; let local_id = vec2(thread_index % GROUP_SIZE.x, thread_index / GROUP_SIZE.x); let global_id = (cluster_id * GROUP_SIZE + local_id) * parameters.grid_scale + cluster_offset; - //TODO: also use the offset return vec2(global_id); } @@ -557,6 +557,14 @@ fn main( if (any(vec2(pixel_coord) >= camera.target_size)) { return; } + if (GROUP_VISUALIZE) + { + var rng = random_init(group_id.y * 1000u + group_id.x, 0u); + let h = random_gen(&rng) * 360.0; + let color = hsv_to_rgb(h, 0.5, 1.0); + textureStore(out_diffuse, pixel_coord, vec4(color, 1.0)); + return; + } let global_index = u32(pixel_coord.y) * camera.target_size.x + u32(pixel_coord.x); var rng = random_init(global_index, parameters.frame_index); diff --git a/blade-render/src/render/mod.rs b/blade-render/src/render/mod.rs index 9e4c03bf..53f8cb97 100644 --- a/blade-render/src/render/mod.rs +++ b/blade-render/src/render/mod.rs @@ -11,7 +11,6 @@ pub use env_map::EnvironmentMap; use std::{collections::HashMap, mem, num::NonZeroU32, path::Path, ptr}; const MAX_RESOURCES: u32 = 1000; -const GRID_SCALES: u32 = 4; const RADIANCE_FORMAT: blade_graphics::TextureFormat = blade_graphics::TextureFormat::Rgba16Float; fn mat4_transform(t: &blade_graphics::Transform) -> glam::Mat4 { @@ -96,9 +95,9 @@ pub struct RayConfig { pub spatial_tap_history: u32, /// Minimal distance to a spatially reused pixel (in the current frame). pub spatial_min_distance: u32, - /// Enable jittering of the compute grid, to allow spatial samples to mix + /// Scale and mix the groups into clusters, to allow spatial samples to mix /// outside of the original workgroup pixel bounds. - pub spatial_jitter: bool, + pub group_mixer: u32, pub t_start: f32, } @@ -300,20 +299,6 @@ struct Blur { atrous_pipeline: blade_graphics::ComputePipeline, } -#[derive(Clone, Copy)] -struct GridMapping { - offset: mint::Vector2, - scale: mint::Vector2, -} -impl Default for GridMapping { - fn default() -> Self { - Self { - offset: [0; 2].into(), - scale: [1; 2].into(), - } - } -} - /// Blade Renderer is a comprehensive rendering solution for /// end user applications. /// @@ -341,7 +326,6 @@ pub struct Renderer { textures: blade_graphics::TextureArray, samplers: Samplers, reservoir_size: u32, - grid_mapping: GridMapping, debug: DebugRender, surface_size: blade_graphics::Extent, surface_info: blade_graphics::SurfaceInfo, @@ -351,7 +335,6 @@ pub struct Renderer { // This way we can embed user info into the allocator. texture_resource_lookup: HashMap>, - random: nanorand::WyRand, } #[repr(C)] @@ -387,7 +370,6 @@ struct MainParams { spatial_min_distance: u32, t_start: f32, use_motion_vectors: u32, - grid_offset: [u32; 2], grid_scale: [u32; 2], } @@ -743,14 +725,12 @@ impl Renderer { textures: blade_graphics::TextureArray::new(), samplers, reservoir_size: sp.reservoir_size, - grid_mapping: GridMapping::default(), debug, surface_size: config.surface_size, surface_info: config.surface_info, frame_index: 0, frame_scene_built: 0, texture_resource_lookup: HashMap::default(), - random: nanorand::WyRand::new(), } } @@ -1128,21 +1108,6 @@ impl Renderer { } self.targets.camera_params[self.frame_index % 2] = self.make_camera_params(camera); self.post_proc_input_index = self.frame_index % 2; - - self.grid_mapping = { - let wg_size = self.main_pipeline.get_workgroup_size(); - let random = nanorand::Rng::generate::(&mut self.random); - let r_offset = random as u32; - let r_scale = (random >> 32) as u32; - GridMapping { - offset: [r_offset % wg_size[0], (r_offset / wg_size[0]) % wg_size[1]].into(), - scale: [ - 1 << (r_scale % GRID_SCALES), - 1 << ((r_scale / GRID_SCALES) % GRID_SCALES), - ] - .into(), - } - }; } /// Ray trace the scene. @@ -1186,11 +1151,12 @@ impl Renderer { } if let mut pass = command_encoder.compute() { - let grid_mapping = if ray_config.spatial_jitter { - self.grid_mapping - } else { - GridMapping::default() + let grid_scale = { + let limit = ray_config.group_mixer; + let r = self.frame_index as u32 ^ 0x5A; + [r % limit + 1, (r / limit) % limit + 1] }; + let groups = self.main_pipeline.get_dispatch_for(self.surface_size); let mut pc = pass.with(&self.main_pipeline); pc.bind( @@ -1211,8 +1177,7 @@ impl Renderer { spatial_min_distance: ray_config.spatial_min_distance, t_start: ray_config.t_start, use_motion_vectors: (self.frame_scene_built == self.frame_index) as u32, - grid_offset: grid_mapping.offset.into(), - grid_scale: grid_mapping.scale.into(), + grid_scale, }, acc_struct: self.acceleration_structure, prev_acc_struct: if self.frame_scene_built < self.frame_index diff --git a/examples/scene/main.rs b/examples/scene/main.rs index 632e9f95..903a30ca 100644 --- a/examples/scene/main.rs +++ b/examples/scene/main.rs @@ -265,7 +265,7 @@ impl Example { spatial_taps: 1, spatial_tap_history: 5, spatial_min_distance: 4, - spatial_jitter: true, + group_mixer: 10, t_start: 0.1, }, denoiser_enabled: true, diff --git a/src/lib.rs b/src/lib.rs index 56d508f8..b7f2b3bb 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -490,7 +490,7 @@ impl Engine { spatial_taps: 1, spatial_tap_history: 5, spatial_min_distance: 4, - spatial_jitter: true, + group_mixer: 10, t_start: 0.01, }, denoiser_enabled: true,