From 0328c633207af84d8517c0fec657d2630c287acc Mon Sep 17 00:00:00 2001 From: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> Date: Thu, 15 Oct 2026 15:03:15 -0500 Subject: [PATCH 1/4] perf(rendering): convert YUV420P frames in 2x2 blocks Signed-off-by: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> --- .../src/shaders/yuv420p_to_rgba.wgsl | 26 ++++++++++++------- crates/rendering/src/yuv_converter.rs | 4 +-- 2 files changed, 19 insertions(+), 11 deletions(-) diff --git a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl index 0d04d30fede..6f756e1bff6 100644 --- a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl +++ b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl @@ -3,18 +3,12 @@ @group(0) @binding(2) var v_plane: texture_2d; @group(0) @binding(3) var output: texture_storage_2d; +// One invocation per 2x2 block: the block shares its chroma sample, so this +// loads U and V once instead of four times (~35% faster on Intel Gen9). @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) global_id: vec3) { - let coords = global_id.xy; + let uv_coords = global_id.xy; let dims = textureDimensions(output); - - if (coords.x >= dims.x || coords.y >= dims.y) { - return; - } - - let y_raw = textureLoad(y_plane, coords, 0).r; - - let uv_coords = coords / 2; let u_dims = textureDimensions(u_plane); let v_dims = textureDimensions(v_plane); let u_clamped = min(uv_coords, u_dims - vec2(1, 1)); @@ -22,6 +16,20 @@ fn main(@builtin(global_invocation_id) global_id: vec3) { let u_raw = textureLoad(u_plane, u_clamped, 0).r; let v_raw = textureLoad(v_plane, v_clamped, 0).r; + let origin = uv_coords * 2u; + convert_pixel(origin, dims, u_raw, v_raw); + convert_pixel(origin + vec2(1u, 0u), dims, u_raw, v_raw); + convert_pixel(origin + vec2(0u, 1u), dims, u_raw, v_raw); + convert_pixel(origin + vec2(1u, 1u), dims, u_raw, v_raw); +} + +fn convert_pixel(coords: vec2, dims: vec2, u_raw: f32, v_raw: f32) { + if (coords.x >= dims.x || coords.y >= dims.y) { + return; + } + + let y_raw = textureLoad(y_plane, coords, 0).r; + let y = (y_raw - 0.0625) * 1.164; let u = u_raw - 0.5; let v = v_raw - 0.5; diff --git a/crates/rendering/src/yuv_converter.rs b/crates/rendering/src/yuv_converter.rs index 4bd420ddb2a..73519a0b9e0 100644 --- a/crates/rendering/src/yuv_converter.rs +++ b/crates/rendering/src/yuv_converter.rs @@ -1168,7 +1168,7 @@ impl YuvToRgbaConverter { }); compute_pass.set_pipeline(&self.pipelines.yuv420p_pipeline); compute_pass.set_bind_group(0, bind_group, &[]); - compute_pass.dispatch_workgroups(width.div_ceil(8), height.div_ceil(8), 1); + compute_pass.dispatch_workgroups(width.div_ceil(16), height.div_ceil(16), 1); } queue.submit(std::iter::once(encoder.finish())); @@ -1239,7 +1239,7 @@ impl YuvToRgbaConverter { }); compute_pass.set_pipeline(&self.pipelines.yuv420p_pipeline); compute_pass.set_bind_group(0, bind_group, &[]); - compute_pass.dispatch_workgroups(width.div_ceil(8), height.div_ceil(8), 1); + compute_pass.dispatch_workgroups(width.div_ceil(16), height.div_ceil(16), 1); } Ok(self.current_output_view()) From 09044e186aa6bcc10c725bca8131bd29ac65c3e6 Mon Sep 17 00:00:00 2001 From: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> Date: Thu, 15 Oct 2026 15:48:49 -0500 Subject: [PATCH 2/4] refactor(style): clean up comment Signed-off-by: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> --- crates/rendering/src/shaders/yuv420p_to_rgba.wgsl | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl index 6f756e1bff6..ddf8a61cc56 100644 --- a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl +++ b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl @@ -3,8 +3,7 @@ @group(0) @binding(2) var v_plane: texture_2d; @group(0) @binding(3) var output: texture_storage_2d; -// One invocation per 2x2 block: the block shares its chroma sample, so this -// loads U and V once instead of four times (~35% faster on Intel Gen9). +// One invocation per 2x2 block: the block shares its chroma sample, so this loads U and V once instead of four times. @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) global_id: vec3) { let uv_coords = global_id.xy; From 6d13723967abe4f2525e346a740890158cc622b2 Mon Sep 17 00:00:00 2001 From: The-Best-Codes <106822363+The-Best-Codes@users.noreply.github.com> Date: Fri, 16 Oct 2026 01:04:26 -0500 Subject: [PATCH 3/4] fix(perf): gate the refactor to only affect linux --- .../src/shaders/yuv420p_to_rgba.wgsl | 33 ++++++++++++++----- crates/rendering/src/yuv_converter.rs | 24 ++++++++++++-- 2 files changed, 45 insertions(+), 12 deletions(-) diff --git a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl index ddf8a61cc56..65f746ab7a5 100644 --- a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl +++ b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl @@ -3,23 +3,38 @@ @group(0) @binding(2) var v_plane: texture_2d; @group(0) @binding(3) var output: texture_storage_2d; -// One invocation per 2x2 block: the block shares its chroma sample, so this loads U and V once instead of four times. +// One invocation per pixel. @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) global_id: vec3) { + let coords = global_id.xy; + let dims = textureDimensions(output); + let uv = load_uv(coords / 2u); + convert_pixel(coords, dims, uv.x, uv.y); +} + +// One invocation per 2x2 block: the block shares its chroma sample, so this loads U and V once instead of four times. +@compute @workgroup_size(8, 8) +fn main_block(@builtin(global_invocation_id) global_id: vec3) { let uv_coords = global_id.xy; let dims = textureDimensions(output); + let uv = load_uv(uv_coords); + + let origin = uv_coords * 2u; + convert_pixel(origin, dims, uv.x, uv.y); + convert_pixel(origin + vec2(1u, 0u), dims, uv.x, uv.y); + convert_pixel(origin + vec2(0u, 1u), dims, uv.x, uv.y); + convert_pixel(origin + vec2(1u, 1u), dims, uv.x, uv.y); +} + +fn load_uv(uv_coords: vec2) -> vec2 { let u_dims = textureDimensions(u_plane); let v_dims = textureDimensions(v_plane); let u_clamped = min(uv_coords, u_dims - vec2(1, 1)); let v_clamped = min(uv_coords, v_dims - vec2(1, 1)); - let u_raw = textureLoad(u_plane, u_clamped, 0).r; - let v_raw = textureLoad(v_plane, v_clamped, 0).r; - - let origin = uv_coords * 2u; - convert_pixel(origin, dims, u_raw, v_raw); - convert_pixel(origin + vec2(1u, 0u), dims, u_raw, v_raw); - convert_pixel(origin + vec2(0u, 1u), dims, u_raw, v_raw); - convert_pixel(origin + vec2(1u, 1u), dims, u_raw, v_raw); + return vec2( + textureLoad(u_plane, u_clamped, 0).r, + textureLoad(v_plane, v_clamped, 0).r + ); } fn convert_pixel(coords: vec2, dims: vec2, u_raw: f32, v_raw: f32) { diff --git a/crates/rendering/src/yuv_converter.rs b/crates/rendering/src/yuv_converter.rs index 73519a0b9e0..167991374da 100644 --- a/crates/rendering/src/yuv_converter.rs +++ b/crates/rendering/src/yuv_converter.rs @@ -107,6 +107,16 @@ const INITIAL_TEXTURE_HEIGHT: u32 = 1080; const TEXTURE_SIZE_PADDING: u32 = 64; +// The 2x2 block kernel only pays off on Linux; on Apple GPUs it measured slower than one invocation per pixel. +#[cfg(target_os = "linux")] +const YUV420P_ENTRY_POINT: &str = "main_block"; +#[cfg(target_os = "linux")] +const YUV420P_PIXELS_PER_WORKGROUP: u32 = 16; +#[cfg(not(target_os = "linux"))] +const YUV420P_ENTRY_POINT: &str = "main"; +#[cfg(not(target_os = "linux"))] +const YUV420P_PIXELS_PER_WORKGROUP: u32 = 8; + fn align_dimension(dim: u32) -> u32 { dim.div_ceil(TEXTURE_SIZE_PADDING) * TEXTURE_SIZE_PADDING } @@ -394,7 +404,7 @@ impl YuvConverterPipelines { label: Some("YUV420P Converter Pipeline"), layout: Some(&yuv420p_pipeline_layout), module: &yuv420p_shader, - entry_point: Some("main"), + entry_point: Some(YUV420P_ENTRY_POINT), compilation_options: Default::default(), cache: None, }); @@ -1168,7 +1178,11 @@ impl YuvToRgbaConverter { }); compute_pass.set_pipeline(&self.pipelines.yuv420p_pipeline); compute_pass.set_bind_group(0, bind_group, &[]); - compute_pass.dispatch_workgroups(width.div_ceil(16), height.div_ceil(16), 1); + compute_pass.dispatch_workgroups( + width.div_ceil(YUV420P_PIXELS_PER_WORKGROUP), + height.div_ceil(YUV420P_PIXELS_PER_WORKGROUP), + 1, + ); } queue.submit(std::iter::once(encoder.finish())); @@ -1239,7 +1253,11 @@ impl YuvToRgbaConverter { }); compute_pass.set_pipeline(&self.pipelines.yuv420p_pipeline); compute_pass.set_bind_group(0, bind_group, &[]); - compute_pass.dispatch_workgroups(width.div_ceil(16), height.div_ceil(16), 1); + compute_pass.dispatch_workgroups( + width.div_ceil(YUV420P_PIXELS_PER_WORKGROUP), + height.div_ceil(YUV420P_PIXELS_PER_WORKGROUP), + 1, + ); } Ok(self.current_output_view()) From 1ae3496473908f66e29c1c9f60b83e55c7d59929 Mon Sep 17 00:00:00 2001 From: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> Date: Fri, 16 Oct 2026 01:29:41 -0500 Subject: [PATCH 4/4] chore(comments): clarify comments to be clearer about what I did Signed-off-by: BestCodes <106822363+The-Best-Codes@users.noreply.github.com> --- crates/rendering/src/shaders/yuv420p_to_rgba.wgsl | 1 - crates/rendering/src/yuv_converter.rs | 2 +- 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl index 65f746ab7a5..5cb3fe842be 100644 --- a/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl +++ b/crates/rendering/src/shaders/yuv420p_to_rgba.wgsl @@ -3,7 +3,6 @@ @group(0) @binding(2) var v_plane: texture_2d; @group(0) @binding(3) var output: texture_storage_2d; -// One invocation per pixel. @compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) global_id: vec3) { let coords = global_id.xy; diff --git a/crates/rendering/src/yuv_converter.rs b/crates/rendering/src/yuv_converter.rs index 167991374da..f0e7fe84809 100644 --- a/crates/rendering/src/yuv_converter.rs +++ b/crates/rendering/src/yuv_converter.rs @@ -107,7 +107,7 @@ const INITIAL_TEXTURE_HEIGHT: u32 = 1080; const TEXTURE_SIZE_PADDING: u32 = 64; -// The 2x2 block kernel only pays off on Linux; on Apple GPUs it measured slower than one invocation per pixel. +// The 2x2 block kernel was only measured as a win on Linux; Apple GPUs were slower and Windows is untested. #[cfg(target_os = "linux")] const YUV420P_ENTRY_POINT: &str = "main_block"; #[cfg(target_os = "linux")]