diff --git a/native/shared/src/renderer/final_pass.rs b/native/shared/src/renderer/final_pass.rs index 4b3753b7..6b1fb942 100644 --- a/native/shared/src/renderer/final_pass.rs +++ b/native/shared/src/renderer/final_pass.rs @@ -107,6 +107,10 @@ impl Renderer { binding: 7, resource: wgpu::BindingResource::TextureView(&self.depth_view), }, + wgpu::BindGroupEntry { + binding: 8, + resource: wgpu::BindingResource::TextureView(&self.tonemap_lut.view), + }, ], }); self.composite_bind_group_cache[composite_cache_index] = Some(composite_bg); diff --git a/native/shared/src/renderer/mod.rs b/native/shared/src/renderer/mod.rs index 137f9ab4..f4d5b664 100644 --- a/native/shared/src/renderer/mod.rs +++ b/native/shared/src/renderer/mod.rs @@ -115,6 +115,7 @@ mod temporal_diagnostics; mod temporal_history; mod temporal_reactive; mod texture_store; +mod tonemap_lut; mod transmitted_shadows; mod transparent_gi; #[cfg(feature = "models3d")] @@ -609,6 +610,8 @@ pub struct Renderer { /// the tonemap. pub composite_pipeline: LazyRenderPipeline, pub composite_layout: wgpu::BindGroupLayout, + /// The AgX tonemap as a 3D table (tonemap_lut.rs), read by the composite. + tonemap_lut: tonemap_lut::TonemapLut, pub composite_sampler: wgpu::Sampler, composite_bind_group_cache: [Option; postfx_chain::CompositeSource::COUNT * 2], /// 0 = ACES (default, matches bloom-reference), 1 = AgX. @@ -4123,6 +4126,7 @@ impl Renderer { label: Some("composite_shader"), source: wgpu::ShaderSource::Wgsl(COMPOSITE_SHADER_WGSL.into()), }); + let tonemap_lut = tonemap_lut::TonemapLut::new(&device, &queue); let composite_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("composite_layout"), entries: &[ @@ -4194,6 +4198,16 @@ impl Renderer { }, count: None, }, + wgpu::BindGroupLayoutEntry { + binding: 8, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D3, + multisampled: false, + }, + count: None, + }, ], }); let composite_pl_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { @@ -8128,6 +8142,7 @@ impl Renderer { scene_compose_bind_group_cache: std::array::from_fn(|_| None), composite_pipeline, composite_layout, + tonemap_lut, composite_sampler, composite_bind_group_cache: std::array::from_fn(|_| None), // 1 = AgX (Troy Sobotka 2022). Matches Blender 4.0+ and diff --git a/native/shared/src/renderer/shaders/mod.rs b/native/shared/src/renderer/shaders/mod.rs index f04d7a0f..17299d9c 100644 --- a/native/shared/src/renderer/shaders/mod.rs +++ b/native/shared/src/renderer/shaders/mod.rs @@ -413,6 +413,7 @@ mod post; pub(super) use post::{ BLOOM_SHADER_WGSL, COMPOSITE_SHADER_WGSL, DOF_SHADER_WGSL, EXPOSURE_SHADER_WGSL, MOTION_BLUR_SHADER_WGSL, SCENE_COMPOSE_SHADER_WGSL, SSS_SHADER_WGSL, TAA_SHADER_WGSL, + TONEMAP_LUT_BAKE_WGSL, }; mod post_upscale; pub(super) use post_upscale::{RCAS_SHADER_WGSL, UPSCALE_SHADER_WGSL}; diff --git a/native/shared/src/renderer/shaders/post.rs b/native/shared/src/renderer/shaders/post.rs index f14d501e..9a995dcd 100644 --- a/native/shared/src/renderer/shaders/post.rs +++ b/native/shared/src/renderer/shaders/post.rs @@ -1520,7 +1520,10 @@ fn fs_main() -> @location(0) vec4 { /// reads hdr_rt and writes ACES-tonemapped linear-RGB. Hardware /// performs the linear→sRGB encode on write because the surface /// format is sRGB. -pub(in crate::renderer) const COMPOSITE_SHADER_WGSL: &str = " +pub(in crate::renderer) const COMPOSITE_SHADER_WGSL: &str = concat!( + include_str!("tonemap_lut.wgsl"), + include_str!("tonemap.wgsl"), + " struct CompositeParams { /// x = tonemap mode (0 = ACES, 1 = AgX) /// y = auto-exposure enabled (0 = off, uses manual x) @@ -1549,6 +1552,7 @@ struct CompositeParams { @group(0) @binding(5) var ssao_tex: texture_2d; @group(0) @binding(6) var ssao_samp: sampler; @group(0) @binding(7) var scene_depth_tex: texture_depth_2d; +@group(0) @binding(8) var tonemap_lut: texture_3d; struct VsOut { @builtin(position) clip_pos: vec4, @@ -1565,106 +1569,14 @@ fn vs_main(@builtin(vertex_index) vid: u32) -> VsOut { return out; } -fn aces_tone(c: vec3) -> vec3 { - let a = 2.51; - let b = 0.03; - let cc = 2.43; - let d = 0.59; - let e = 0.14; - return clamp((c * (c * a + b)) / (c * (c * cc + d) + e), vec3(0.0), vec3(1.0)); -} - -// ACES RRT+ODT fit with an explicit 6.0 white point. This is the MJP -// BakingLab fit used by Godot 4's ACES mode, including its 1.8 exposure bias -// and gamut transforms. Keep the compact Narkowicz fit above as mode 0 for -// output compatibility; mode 2 is the reference-quality/highlight-preserving -// path used by scenes that request the full transform. -fn aces_full_tone(c_in: vec3) -> vec3 { - let exposure_bias = 1.8; - let a = 0.0245786; - let b = 0.000090537; - let cc = 0.983729; - let d = 0.432951; - let e = 0.238081; - let rgb_to_rrt = mat3x3( - vec3(0.59719, 0.35458, 0.04823) * exposure_bias, - vec3(0.07600, 0.90834, 0.01566) * exposure_bias, - vec3(0.02840, 0.13383, 0.83777) * exposure_bias, - ); - let odt_to_rgb = mat3x3( - vec3( 1.60475, -0.53108, -0.07367), - vec3(-0.10208, 1.10813, -0.00605), - vec3(-0.00327, -0.07276, 1.07602), - ); - // Godot's source uses row-vector multiplication (`color *= matrix`). - // Preserve that convention explicitly; swapping this order transposes the - // gamut transform and gives neutral surfaces a strong green cast. - let c = max(c_in, vec3(0.0)) * rgb_to_rrt; - let mapped = (c * (c + a) - b) / (c * (cc * c + d) + e); - let white = 6.0 * exposure_bias; - let white_mapped = (white * (white + a) - b) - / (white * (cc * white + d) + e); - return clamp(mapped * odt_to_rgb / white_mapped, - vec3(0.0), vec3(1.0)); -} - -// --- AgX tonemap (Blender/Filament reference) --- -// Better hue preservation than ACES in saturated regions — reds -// stay red instead of shifting toward orange, blues stay blue -// instead of shifting toward cyan. Same sigmoid shape overall, -// so the overall contrast is similar. - -fn agx_default_contrast_approx(x: vec3) -> vec3 { - let x2 = x * x; - let x4 = x2 * x2; - return vec3(15.5) * x4 * x2 - - vec3(40.14) * x4 * x - + vec3(31.96) * x4 - - vec3(6.868) * x2 * x - + vec3(0.4298) * x2 - + vec3(0.1191) * x - - vec3(0.00232); -} - -fn agx_tone(val_in: vec3) -> vec3 { - // AgX input transform — compresses the input color gamut. - let agx_mat = mat3x3( - vec3(0.842479062253094, 0.0423282422610123, 0.0423756549057051), - vec3(0.0784335999999992, 0.878468636469772, 0.0784336), - vec3(0.0792237451477643, 0.0791661274605434, 0.879142973793104), - ); - // Log2-space normalization range. Anything outside gets clamped - // — the sigmoid maps this window to [0, 1]. - let min_ev = -12.47393; - let max_ev = 4.026069; - - var val = agx_mat * val_in; - // Log2 encode, clamp to range, normalize to [0, 1]. - val = max(val, vec3(1e-10)); - val = clamp(log2(val), vec3(min_ev), vec3(max_ev)); - val = (val - vec3(min_ev)) / (max_ev - min_ev); - - // Sigmoid contrast curve. - val = agx_default_contrast_approx(val); - return val; -} - -fn agx_eotf(val_in: vec3) -> vec3 { - // AgX inverse input transform — re-expands back to target - // display gamut. The surface is sRGB-format so hardware - // applies the sRGB EOTF on write; we output linear here. - let agx_mat_inv = mat3x3( - vec3( 1.19687900512017, -0.0528968517574562, -0.0529716355144438), - vec3(-0.0980208811401368, 1.15190312990417, -0.0980434501171241), - vec3(-0.0990297440797205, -0.0989611768448433, 1.15107367264116), - ); - return agx_mat_inv * val_in; -} - // Tonemap — branches on u.params.x (0 = compact ACES, 1 = AgX, -// 2 = full ACES RRT+ODT fit). Extracted -// so the sharpen pass can tonemap neighbour HDR samples through the -// same path as the center pixel. +// 2 = full ACES RRT+ODT fit; tonemap.wgsl). Extracted so the sharpen pass can +// tonemap neighbour HDR samples through the same path as the center pixel. +// AgX is read from its baked table (renderer/tonemap_lut.rs): the curve costs +// dozens of ALU ops and fifteen transcendentals, and this runs up to four times +// per pixel (centre, sharpen average, chromatic aberration R and B). On mobile +// GPUs that dominated the pass: the live state of two AgX evaluations cost a +// Pixel 4a 21 ms of a 25 ms composite at 1080x2340. fn tonemap_select(hdr: vec3) -> vec3 { if (u.params.x < 0.5) { return aces_tone(hdr); @@ -1672,43 +1584,7 @@ fn tonemap_select(hdr: vec3) -> vec3 { if (u.params.x > 1.5) { return aces_full_tone(hdr); } - // Full Filament AgX pipeline: - // agx_tone (inset + log2 + sigmoid) - // agx_eotf (outset = inverse inset) - // pow(2.2) (display-encoded → linear, required for sRGB surface) - // - // Filament's polynomial-approximation AgX (what we use) systematically - // under-saturates vs Blender/Cycles' LUT-based AgX. Apply the 'Punchy' - // look — saturation 1.4, slight contrast — as a post step to bring - // colours closer to the Cycles ground-truth reference. Matches - // Filament::ToneMapper::AGX_PUNCHY. - let hdr_safe = max(hdr, vec3(0.0)); - let agx_display = clamp(agx_eotf(agx_tone(hdr_safe)), vec3(0.0), vec3(1.0)); - let linear = pow(agx_display, vec3(2.2)); - // Punchy look: post-tonemap saturation + contrast in display space. - // slope/offset/power are per-channel lift/gamma/gain; saturation - // multiplies chroma around the pixel's luma. - let agx_punchy = agx_look_punchy(linear); - return agx_punchy; -} - -fn agx_look_punchy(val: vec3) -> vec3 { - // Per-channel slope (gain): subtle warm white-balance toward the - // 'golden hour' feel of Cycles-AgX renders — +3% red, -2% blue. - // The Bistro outdoor.hdr leans cool; without this, the shadows - // and IBL fill pull the overall image toward blue-grey vs. - // Cycles's warmer midtone neutral. - let slope = vec3(1.03, 1.00, 0.98); - let offset = vec3(0.0); - // A whisker above neutral — Cycles-AgX with its LUT-based DRT - // preserves enough chroma on its own that our polynomial fit only - // needs a gentle post-boost to catch up. - let power = vec3(1.1); - let saturation = 1.1; - // ASC-CDL-ish: (val * slope + offset) ^ power - let toned = pow(max(val * slope + offset, vec3(0.0)), power); - let luma = dot(toned, vec3(0.2126, 0.7152, 0.0722)); - return vec3(luma) + saturation * (toned - vec3(luma)); + return textureSampleLevel(tonemap_lut, hdr_samp, tonemap_lut_uvw(hdr), 0.0).rgb; } // Hash-based pseudo-random in [0, 1). Cheap noise function for grain; @@ -2005,4 +1881,23 @@ fn fs_main(in: VsOut) -> @location(0) vec4 { return vec4(ldr, 1.0); } -"; +", +); + +/// Bakes agx_curve() into the AgX table (renderer/tonemap_lut.rs): one texel +/// per HDR colour on the table's grid. +pub(in crate::renderer) const TONEMAP_LUT_BAKE_WGSL: &str = concat!( + include_str!("tonemap_lut.wgsl"), + include_str!("tonemap.wgsl"), + " +@group(0) @binding(0) var lut: texture_storage_3d; + +@compute @workgroup_size(4, 4, 4) +fn bake(@builtin(global_invocation_id) id: vec3) { + if (any(id >= vec3(u32(TONEMAP_LUT_SIZE)))) { + return; + } + textureStore(lut, id, vec4(agx_curve(tonemap_lut_hdr(id)), 1.0)); +} +", +); diff --git a/native/shared/src/renderer/shaders/tonemap.wgsl b/native/shared/src/renderer/shaders/tonemap.wgsl new file mode 100644 index 00000000..fcb1f44d --- /dev/null +++ b/native/shared/src/renderer/shaders/tonemap.wgsl @@ -0,0 +1,141 @@ +// Bloom's display tonemap curves: compact ACES (kind 0), AgX (kind 1) and the +// full ACES fit (kind 2). The ACES curves are cheap and evaluated per pixel; +// AgX is baked into a table (renderer/tonemap_lut.rs, tonemap_lut.wgsl). + +fn aces_tone(c: vec3) -> vec3 { + let a = 2.51; + let b = 0.03; + let cc = 2.43; + let d = 0.59; + let e = 0.14; + return clamp((c * (c * a + b)) / (c * (c * cc + d) + e), vec3(0.0), vec3(1.0)); +} + +// ACES RRT+ODT fit with an explicit 6.0 white point. This is the MJP +// BakingLab fit used by Godot 4's ACES mode, including its 1.8 exposure bias +// and gamut transforms. Keep the compact Narkowicz fit above as mode 0 for +// output compatibility; mode 2 is the reference-quality/highlight-preserving +// path used by scenes that request the full transform. +fn aces_full_tone(c_in: vec3) -> vec3 { + let exposure_bias = 1.8; + let a = 0.0245786; + let b = 0.000090537; + let cc = 0.983729; + let d = 0.432951; + let e = 0.238081; + let rgb_to_rrt = mat3x3( + vec3(0.59719, 0.35458, 0.04823) * exposure_bias, + vec3(0.07600, 0.90834, 0.01566) * exposure_bias, + vec3(0.02840, 0.13383, 0.83777) * exposure_bias, + ); + let odt_to_rgb = mat3x3( + vec3( 1.60475, -0.53108, -0.07367), + vec3(-0.10208, 1.10813, -0.00605), + vec3(-0.00327, -0.07276, 1.07602), + ); + // Godot's source uses row-vector multiplication (`color *= matrix`). + // Preserve that convention explicitly; swapping this order transposes the + // gamut transform and gives neutral surfaces a strong green cast. + let c = max(c_in, vec3(0.0)) * rgb_to_rrt; + let mapped = (c * (c + a) - b) / (c * (cc * c + d) + e); + let white = 6.0 * exposure_bias; + let white_mapped = (white * (white + a) - b) + / (white * (cc * white + d) + e); + return clamp(mapped * odt_to_rgb / white_mapped, + vec3(0.0), vec3(1.0)); +} + +// --- AgX tonemap (Blender/Filament reference) --- +// Better hue preservation than ACES in saturated regions — reds +// stay red instead of shifting toward orange, blues stay blue +// instead of shifting toward cyan. Same sigmoid shape overall, +// so the overall contrast is similar. + +fn agx_default_contrast_approx(x: vec3) -> vec3 { + let x2 = x * x; + let x4 = x2 * x2; + return vec3(15.5) * x4 * x2 + - vec3(40.14) * x4 * x + + vec3(31.96) * x4 + - vec3(6.868) * x2 * x + + vec3(0.4298) * x2 + + vec3(0.1191) * x + - vec3(0.00232); +} + +fn agx_tone(val_in: vec3) -> vec3 { + // AgX input transform — compresses the input color gamut. + let agx_mat = mat3x3( + vec3(0.842479062253094, 0.0423282422610123, 0.0423756549057051), + vec3(0.0784335999999992, 0.878468636469772, 0.0784336), + vec3(0.0792237451477643, 0.0791661274605434, 0.879142973793104), + ); + // Log2-space normalization range. Anything outside gets clamped + // — the sigmoid maps this window to [0, 1]. + let min_ev = -12.47393; + let max_ev = 4.026069; + + var val = agx_mat * val_in; + // Log2 encode, clamp to range, normalize to [0, 1]. + val = max(val, vec3(1e-10)); + val = clamp(log2(val), vec3(min_ev), vec3(max_ev)); + val = (val - vec3(min_ev)) / (max_ev - min_ev); + + // Sigmoid contrast curve. + val = agx_default_contrast_approx(val); + return val; +} + +fn agx_eotf(val_in: vec3) -> vec3 { + // AgX inverse input transform — re-expands back to target + // display gamut. The surface is sRGB-format so hardware + // applies the sRGB EOTF on write; we output linear here. + let agx_mat_inv = mat3x3( + vec3( 1.19687900512017, -0.0528968517574562, -0.0529716355144438), + vec3(-0.0980208811401368, 1.15190312990417, -0.0980434501171241), + vec3(-0.0990297440797205, -0.0989611768448433, 1.15107367264116), + ); + return agx_mat_inv * val_in; +} + +// The AgX display curve. The composite pass reads it from the table baked +// from this function (renderer/tonemap_lut.rs) rather than evaluating it. +fn agx_curve(hdr: vec3) -> vec3 { + // Full Filament AgX pipeline: + // agx_tone (inset + log2 + sigmoid) + // agx_eotf (outset = inverse inset) + // pow(2.2) (display-encoded → linear, required for sRGB surface) + // + // Filament's polynomial-approximation AgX (what we use) systematically + // under-saturates vs Blender/Cycles' LUT-based AgX. Apply the 'Punchy' + // look — saturation 1.4, slight contrast — as a post step to bring + // colours closer to the Cycles ground-truth reference. Matches + // Filament::ToneMapper::AGX_PUNCHY. + let hdr_safe = max(hdr, vec3(0.0)); + let agx_display = clamp(agx_eotf(agx_tone(hdr_safe)), vec3(0.0), vec3(1.0)); + let linear = pow(agx_display, vec3(2.2)); + // Punchy look: post-tonemap saturation + contrast in display space. + // slope/offset/power are per-channel lift/gamma/gain; saturation + // multiplies chroma around the pixel's luma. + let agx_punchy = agx_look_punchy(linear); + return agx_punchy; +} + +fn agx_look_punchy(val: vec3) -> vec3 { + // Per-channel slope (gain): subtle warm white-balance toward the + // 'golden hour' feel of Cycles-AgX renders — +3% red, -2% blue. + // The Bistro outdoor.hdr leans cool; without this, the shadows + // and IBL fill pull the overall image toward blue-grey vs. + // Cycles's warmer midtone neutral. + let slope = vec3(1.03, 1.00, 0.98); + let offset = vec3(0.0); + // A whisker above neutral — Cycles-AgX with its LUT-based DRT + // preserves enough chroma on its own that our polynomial fit only + // needs a gentle post-boost to catch up. + let power = vec3(1.1); + let saturation = 1.1; + // ASC-CDL-ish: (val * slope + offset) ^ power + let toned = pow(max(val * slope + offset, vec3(0.0)), power); + let luma = dot(toned, vec3(0.2126, 0.7152, 0.0722)); + return vec3(luma) + saturation * (toned - vec3(luma)); +} diff --git a/native/shared/src/renderer/shaders/tonemap_lut.wgsl b/native/shared/src/renderer/shaders/tonemap_lut.wgsl new file mode 100644 index 00000000..bdfb4da8 --- /dev/null +++ b/native/shared/src/renderer/shaders/tonemap_lut.wgsl @@ -0,0 +1,26 @@ +// The AgX curve as a 3D table (renderer/tonemap_lut.rs): the axes are the HDR +// colour's channels in log2, over the range below, one texel at each end. +// Below the range AgX is black (it clamps at 2^-12.47); above it, it is white +// (its inset matrix saturates every channel once any exceeds about 2^8.6). +// Trilinear filtering stays within 1/255 of the curve in mid-tones and 2/255 +// anywhere, display-encoded (tonemap_lut.rs tests). +const TONEMAP_LUT_SIZE: f32 = 64.0; +const TONEMAP_LUT_LOG2_MIN: f32 = -12.5; +const TONEMAP_LUT_LOG2_MAX: f32 = 9.0; + +// Texture coordinate of an HDR colour. +fn tonemap_lut_uvw(hdr: vec3) -> vec3 { + let s = clamp( + (log2(max(hdr, vec3(1e-30))) - TONEMAP_LUT_LOG2_MIN) / + (TONEMAP_LUT_LOG2_MAX - TONEMAP_LUT_LOG2_MIN), + vec3(0.0), + vec3(1.0), + ); + return (s * (TONEMAP_LUT_SIZE - 1.0) + 0.5) / TONEMAP_LUT_SIZE; +} + +// The HDR colour a texel holds. +fn tonemap_lut_hdr(texel: vec3) -> vec3 { + let s = vec3(texel) / (TONEMAP_LUT_SIZE - 1.0); + return exp2(mix(vec3(TONEMAP_LUT_LOG2_MIN), vec3(TONEMAP_LUT_LOG2_MAX), s)); +} diff --git a/native/shared/src/renderer/tonemap_lut.rs b/native/shared/src/renderer/tonemap_lut.rs new file mode 100644 index 00000000..9949b7c2 --- /dev/null +++ b/native/shared/src/renderer/tonemap_lut.rs @@ -0,0 +1,240 @@ +//! The AgX display curve as a 3D table, so the composite pass reads it with +//! one filtered lookup instead of evaluating it per pixel (and up to four times +//! per pixel: centre, sharpen average, chromatic aberration R and B). +//! +//! Baked once on the GPU from the same WGSL curve (`agx_curve()` in +//! `shaders/tonemap.wgsl`); its indexing is in `shaders/tonemap_lut.wgsl`, +//! shared by the bake and the composite. The ACES curves are cheap and stay +//! analytic. + +use super::shaders::TONEMAP_LUT_BAKE_WGSL; + +/// Texels per axis; `TONEMAP_LUT_SIZE` in `shaders/tonemap_lut.wgsl`. +pub(super) const TONEMAP_LUT_SIZE: u32 = 64; +const TONEMAP_LUT_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba16Float; + +pub(super) struct TonemapLut { + _texture: wgpu::Texture, + pub(super) view: wgpu::TextureView, +} + +impl TonemapLut { + /// Create and bake the table (submitted on `queue`). + pub(super) fn new(device: &wgpu::Device, queue: &wgpu::Queue) -> Self { + let texture = device.create_texture(&wgpu::TextureDescriptor { + label: Some("tonemap_lut"), + size: wgpu::Extent3d { + width: TONEMAP_LUT_SIZE, + height: TONEMAP_LUT_SIZE, + depth_or_array_layers: TONEMAP_LUT_SIZE, + }, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D3, + format: TONEMAP_LUT_FORMAT, + usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING, + view_formats: &[], + }); + let view = texture.create_view(&wgpu::TextureViewDescriptor::default()); + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("tonemap_lut_bake"), + source: wgpu::ShaderSource::Wgsl(TONEMAP_LUT_BAKE_WGSL.into()), + }); + let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("tonemap_lut_bake_layout"), + entries: &[wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::WriteOnly, + format: TONEMAP_LUT_FORMAT, + view_dimension: wgpu::TextureViewDimension::D3, + }, + count: None, + }], + }); + let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("tonemap_lut_bake_pl_layout"), + bind_group_layouts: &[Some(&layout)], + immediate_size: 0, + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("tonemap_lut_bake"), + layout: Some(&pipeline_layout), + module: &shader, + entry_point: Some("bake"), + compilation_options: Default::default(), + cache: None, + }); + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("tonemap_lut_bake_bg"), + layout: &layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(&view), + }], + }); + let mut encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("tonemap_lut_bake"), + }); + { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("tonemap_lut_bake"), + timestamp_writes: None, + }); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &bind_group, &[]); + let groups = TONEMAP_LUT_SIZE.div_ceil(4); + pass.dispatch_workgroups(groups, groups, groups); + } + queue.submit([encoder.finish()]); + Self { + _texture: texture, + view, + } + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn wgsl_table_size_matches() { + let wgsl = include_str!("shaders/tonemap_lut.wgsl"); + assert!(wgsl.contains(&format!( + "const TONEMAP_LUT_SIZE: f32 = {TONEMAP_LUT_SIZE}.0;" + ))); + } + + fn device() -> Option<(wgpu::Device, wgpu::Queue)> { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor { + backends: wgpu::Backends::all(), + ..wgpu::InstanceDescriptor::new_without_display_handle() + }); + let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions { + power_preference: wgpu::PowerPreference::LowPower, + compatible_surface: None, + force_fallback_adapter: false, + })) + .ok()?; + pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + label: Some("tonemap-lut-test-device"), + ..Default::default() + })) + .ok() + } + + /// Display-encoded (sRGB) difference between the lookup the composite pass + /// does and agx_curve(), over HDR colours from 2^-14 to 2^10 in every + /// combination of channels. Skipped without a GPU. + #[test] + fn table_matches_agx_curve() { + let Some((device, queue)) = device() else { + return; + }; + const N: u32 = 64; + let check = format!( + "{}{} +@group(0) @binding(0) var lut_tex: texture_3d; +@group(0) @binding(1) var lut_samp: sampler; +@group(0) @binding(2) var worst: array, 2>; + +fn srgb(c: vec3) -> vec3 {{ + let x = clamp(c, vec3(0.0), vec3(1.0)); + return select(1.055 * pow(x, vec3(1.0 / 2.4)) - 0.055, x * 12.92, x <= vec3(0.0031308)); +}} + +@compute @workgroup_size(4, 4, 4) +fn check(@builtin(global_invocation_id) id: vec3) {{ + let hdr = exp2(mix(vec3(-14.0), vec3(10.0), (vec3(id) + 0.37) / {N}.0)); + let exact = srgb(agx_curve(hdr)); + let table = srgb(textureSampleLevel(lut_tex, lut_samp, tonemap_lut_uvw(hdr), 0.0).rgb); + let d = abs(exact - table); + let e = u32(max(d.x, max(d.y, d.z)) * 1000000.0); + atomicMax(&worst[0], e); + if (all(hdr > vec3(0.01)) && all(hdr < vec3(4.0))) {{ + atomicMax(&worst[1], e); + }} +}} +", + include_str!("shaders/tonemap_lut.wgsl"), + include_str!("shaders/tonemap.wgsl"), + ); + let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("tonemap_lut_check"), + source: wgpu::ShaderSource::Wgsl(check.into()), + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("tonemap_lut_check"), + layout: None, + module: &module, + entry_point: Some("check"), + compilation_options: Default::default(), + cache: None, + }); + // The composite's sampler: linear, clamped. + let sampler = device.create_sampler(&wgpu::SamplerDescriptor { + address_mode_u: wgpu::AddressMode::ClampToEdge, + address_mode_v: wgpu::AddressMode::ClampToEdge, + address_mode_w: wgpu::AddressMode::ClampToEdge, + mag_filter: wgpu::FilterMode::Linear, + min_filter: wgpu::FilterMode::Linear, + ..Default::default() + }); + let lut = TonemapLut::new(&device, &queue); + let worst = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: 8, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let readback = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: 8, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &pipeline.get_bind_group_layout(0), + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(&lut.view), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::Sampler(&sampler), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: worst.as_entire_binding(), + }, + ], + }); + let mut encoder = device.create_command_encoder(&Default::default()); + { + let mut pass = encoder.begin_compute_pass(&Default::default()); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &bind_group, &[]); + pass.dispatch_workgroups(N / 4, N / 4, N / 4); + } + encoder.copy_buffer_to_buffer(&worst, 0, &readback, 0, 8); + queue.submit([encoder.finish()]); + readback.slice(..).map_async(wgpu::MapMode::Read, |_| {}); + let _ = device.poll(wgpu::PollType::Wait { + submission_index: None, + timeout: None, + }); + let v: Vec = bytemuck::cast_slice(&readback.slice(..).get_mapped_range()).to_vec(); + let (all, mid) = (v[0] as f32 / 1e6, v[1] as f32 / 1e6); + println!( + "AgX table: worst {:.2}/255, mid-tones {:.2}/255", + all * 255.0, + mid * 255.0 + ); + assert!(all < 2.0 / 255.0, "worst {all}"); + assert!(mid < 1.0 / 255.0, "mid-tones {mid}"); + } +}