From 652e5ea63353f00835b020cbf7c86641efd50a42 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Mon, 9 Nov 2026 10:01:20 +0100 Subject: [PATCH 1/2] composite: read AgX from a baked 3D table The composite pass evaluated the AgX curve up to four times per pixel (centre, sharpen average, chromatic aberration R and B): an inset matrix, three log2, a sixth-order polynomial, an outset matrix and two per-channel pow. On mobile GPUs the live state of two evaluations cut occupancy: on a Pixel 4a (Adreno 618) at 1080x2340 one AgX call cost the pass 4 ms, two plus sharpen 25 ms. agx_curve() (now in shaders/tonemap.wgsl with the ACES curves) is baked once into a 64^3 Rgba16Float table indexed by log2 of the HDR colour, 2^-12.5..2^9: AgX clamps below that, and saturates every channel above it. The composite reads it with one filtered lookup through its linear sampler (binding 8). The ACES curves are cheap and stay analytic. The table is within 0.84/255 of the curve in mid-tones and 1.71/255 anywhere, display-encoded (tonemap_lut tests, on any GPU adapter). Pixel 4a, medium: final_composite_pass 25 -> 9 ms. --- native/shared/src/renderer/final_pass.rs | 4 + native/shared/src/renderer/mod.rs | 15 ++ native/shared/src/renderer/shaders/mod.rs | 1 + native/shared/src/renderer/shaders/post.rs | 171 +++----------- .../shared/src/renderer/shaders/tonemap.wgsl | 141 +++++++++++ .../src/renderer/shaders/tonemap_lut.wgsl | 26 +++ native/shared/src/renderer/tonemap_lut.rs | 219 ++++++++++++++++++ 7 files changed, 439 insertions(+), 138 deletions(-) create mode 100644 native/shared/src/renderer/shaders/tonemap.wgsl create mode 100644 native/shared/src/renderer/shaders/tonemap_lut.wgsl create mode 100644 native/shared/src/renderer/tonemap_lut.rs 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..7bc518ba --- /dev/null +++ b/native/shared/src/renderer/tonemap_lut.rs @@ -0,0 +1,219 @@ +//! 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}"); + } +} From c0bafece1bb95ff5c1f2ce646729424c32c2030c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ralph=20K=C3=BCpper?= Date: Mon, 9 Nov 2026 11:41:33 +0100 Subject: [PATCH 2/2] tonemap_lut: rustfmt --- native/shared/src/renderer/tonemap_lut.rs | 35 ++++++++++++++++++----- 1 file changed, 28 insertions(+), 7 deletions(-) diff --git a/native/shared/src/renderer/tonemap_lut.rs b/native/shared/src/renderer/tonemap_lut.rs index 7bc518ba..9949b7c2 100644 --- a/native/shared/src/renderer/tonemap_lut.rs +++ b/native/shared/src/renderer/tonemap_lut.rs @@ -88,7 +88,10 @@ impl TonemapLut { pass.dispatch_workgroups(groups, groups, groups); } queue.submit([encoder.finish()]); - Self { _texture: texture, view } + Self { + _texture: texture, + view, + } } } @@ -99,7 +102,9 @@ mod tests { #[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;"))); + assert!(wgsl.contains(&format!( + "const TONEMAP_LUT_SIZE: f32 = {TONEMAP_LUT_SIZE}.0;" + ))); } fn device() -> Option<(wgpu::Device, wgpu::Queue)> { @@ -194,9 +199,18 @@ fn check(@builtin(global_invocation_id) id: vec3) {{ 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() }, + 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()); @@ -209,10 +223,17 @@ fn check(@builtin(global_invocation_id) id: vec3) {{ 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 _ = 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); + 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}"); }