From ec5ddcf55e293cd3bb4b3b3e9cfa189a97ac630f Mon Sep 17 00:00:00 2001 From: John Tur Date: Mon, 3 Aug 2026 10:37:21 -0400 Subject: [PATCH] Tweak renderer resource management (#62069) Factoring some changes out of https://github.com/zed-industries/zed/pull/61461: - On Metal and wgpu, bulk-upload instance buffers at the start of the frame rather than uploading chunks as we draw them. The depth buffer patch needs this architecture since all opaque quads in the scene are painted before anything else. This also lets us get rid of the awkward retry loop when the GPU buffer is too small for the scene. - On Windows, use a constant buffer for the instance offset rather than creating SRVs with different `FirstElement`s. `FirstElement` is buggy on some vendors' drivers, and reusing the SRV object should reduce D3D runtime overhead. - Misc. refactoring: remove redundant parameters, deduplicate code, create fewer bind groups on wgpu, etc. Fixes https://github.com/zed-industries/zed/issues/49588 Fixes https://github.com/zed-industries/zed/issues/59192 Release Notes: - Fixed an issue where UI elements would flicker on some Windows devices. --- crates/gpui/src/window.rs | 64 +- crates/gpui_macos/src/metal_renderer.rs | 1007 ++++++++----------- crates/gpui_wgpu/src/shaders.wgsl | 4 +- crates/gpui_wgpu/src/wgpu_renderer.rs | 747 +++++++------- crates/gpui_windows/src/directx_renderer.rs | 236 +++-- crates/gpui_windows/src/shaders.hlsl | 24 +- 6 files changed, 951 insertions(+), 1131 deletions(-) diff --git a/crates/gpui/src/window.rs b/crates/gpui/src/window.rs index 3e8563c01b..1b5011e076 100644 --- a/crates/gpui/src/window.rs +++ b/crates/gpui/src/window.rs @@ -3967,6 +3967,45 @@ impl Window { } } + fn largest_border_interior(quad: &Quad) -> Bounds { + let radii = &quad.corner_radii; + let widths = &quad.border_widths; + let edge_radii = Edges { + top: radii.top_left.max(radii.top_right), + right: radii.top_right.max(radii.bottom_right), + bottom: radii.bottom_left.max(radii.bottom_right), + left: radii.top_left.max(radii.bottom_left), + }; + + let antialias_inset = point(ScaledPixels(1.0), ScaledPixels(1.0)); + let inset_bounds = |top_left_inset, bottom_right_inset| { + Bounds::from_corners( + quad.bounds.origin + top_left_inset + antialias_inset, + quad.bounds.bottom_right() - bottom_right_inset - antialias_inset, + ) + }; + + // Rounded corners need only be excluded on one axis. Either candidate + // is empty of border pixels, so use the larger interior. + let horizontal_band = inset_bounds( + point(widths.left, widths.top.max(edge_radii.top)), + point(widths.right, widths.bottom.max(edge_radii.bottom)), + ); + let vertical_band = inset_bounds( + point(widths.left.max(edge_radii.left), widths.top), + point(widths.right.max(edge_radii.right), widths.bottom), + ); + + let area = |bounds: &Bounds| { + bounds.size.width.0.max(0.) * bounds.size.height.0.max(0.) + }; + if area(&horizontal_band) >= area(&vertical_band) { + horizontal_band + } else { + vertical_band + } + } + /// Paint one or more quads into the scene for the next frame at the current stacking context. /// Quads are colored rectangular regions with an optional background, border, and corner radius. /// see [`fill`], [`outline`], and [`quad`] to construct this type. @@ -3998,29 +4037,10 @@ impl Window { return; } - // We're drawing a quad with a border but no fill color. Painting this quad would run the quad shader for every - // transparent interior pixel, which is especially costly when the quad is large. - // Instead, split it into four non-overlapping strips that cover the regions where borders are painted: - // the side strips own the straight left and right edges, while the top and bottom strips own the horizontal - // edges and the rounded corners. - let radii = &quad.corner_radii; - let widths = &quad.border_widths; - - let antialias_slack = point(ScaledPixels(1.0), ScaledPixels(1.0)); - let top_left_inset = point( - widths.left, - widths.top.max(radii.top_left).max(radii.top_right), - ) + antialias_slack; - let bottom_right_inset = point( - widths.right, - widths.bottom.max(radii.bottom_left).max(radii.bottom_right), - ) + antialias_slack; - + // Splitting a border-only quad around its empty interior avoids shading + // every transparent pixel inside large outlines. let outer_bounds = quad.bounds; - let inner_bounds = Bounds::from_corners( - outer_bounds.origin + top_left_inset, - outer_bounds.bottom_right() - bottom_right_inset, - ); + let inner_bounds = Self::largest_border_interior(&quad); if inner_bounds.is_empty() { self.next_frame.scene.insert_primitive(quad); diff --git a/crates/gpui_macos/src/metal_renderer.rs b/crates/gpui_macos/src/metal_renderer.rs index 5a72e23e14..4f6b0fee78 100644 --- a/crates/gpui_macos/src/metal_renderer.rs +++ b/crates/gpui_macos/src/metal_renderer.rs @@ -1,5 +1,5 @@ use crate::metal_atlas::MetalAtlas; -use anyhow::Result; +use anyhow::{Context as _, Result}; use block::ConcreteBlock; use cocoa::{ base::{NO, YES}, @@ -7,9 +7,8 @@ use cocoa::{ quartzcore::AutoresizingMask, }; use gpui::{ - AtlasTextureId, Background, Bounds, ContentMask, DevicePixels, MonochromeSprite, PaintSurface, - Path, Point, PolychromeSprite, PrimitiveBatch, Quad, ScaledPixels, Scene, Shadow, Size, - Surface, Underline, point, size, + AtlasTextureId, Background, Bounds, ContentMask, DevicePixels, PaintSurface, Path, Point, + PrimitiveBatch, ScaledPixels, Scene, Size, point, size, }; #[cfg(any(test, feature = "test-support"))] use image::RgbaImage; @@ -22,12 +21,11 @@ use core_video::{ use foreign_types::{ForeignType, ForeignTypeRef}; use metal::{ CAMetalLayer, CommandQueue, MTLGPUFamily, MTLPixelFormat, MTLResourceOptions, NSRange, - RenderPassColorAttachmentDescriptorRef, }; use objc::{self, msg_send, sel, sel_impl}; use parking_lot::Mutex; -use std::{cell::Cell, ffi::c_void, mem, ptr, sync::Arc}; +use std::{cell::Cell, ffi::c_void, mem, mem::MaybeUninit, ops::Range, ptr, slice, sync::Arc}; // Exported to metal pub(crate) type PointF = gpui::Point; @@ -39,6 +37,9 @@ const SHADERS_SOURCE_FILE: &str = include_str!(concat!(env!("OUT_DIR"), "/stitch // Use 4x MSAA, all devices support it. // https://developer.apple.com/documentation/metal/mtldevice/1433355-supportstexturesamplecount const PATH_SAMPLE_COUNT: u32 = 4; +/// Metal requires the offset a buffer is bound at to be 256-byte aligned. +const INSTANCE_BUFFER_ALIGNMENT: usize = 256; +const MAX_INSTANCE_BUFFER_SIZE: usize = 256 * 1024 * 1024; pub(crate) type Context = Arc>; pub(crate) type Renderer = MetalRenderer; @@ -468,58 +469,68 @@ impl MetalRenderer { return; }; - loop { - let mut instance_buffer = self - .instance_buffer_pool - .lock() - .acquire(&self.device, self.is_unified_memory); - - let command_buffer = - self.draw_primitives(scene, &mut instance_buffer, drawable, viewport_size); - - match command_buffer { - Ok(command_buffer) => { - let instance_buffer_pool = self.instance_buffer_pool.clone(); - let instance_buffer = Cell::new(Some(instance_buffer)); - let block = ConcreteBlock::new(move |_| { - if let Some(instance_buffer) = instance_buffer.take() { - instance_buffer_pool.lock().release(instance_buffer); - } - }); - let block = block.copy(); - command_buffer.add_completed_handler(&block); - - if self.presents_with_transaction { - command_buffer.commit(); - command_buffer.wait_until_scheduled(); - drawable.present(); - } else { - command_buffer.present_drawable(drawable); - command_buffer.commit(); - } - return; - } - Err(err) => { - log::error!( - "failed to render: {}. retrying with larger instance buffer size", - err - ); - let mut instance_buffer_pool = self.instance_buffer_pool.lock(); - let buffer_size = instance_buffer_pool.buffer_size; - if buffer_size >= 256 * 1024 * 1024 { - log::error!("instance buffer size grew too large: {}", buffer_size); - break; - } - instance_buffer_pool.reset(buffer_size * 2); - log::info!( - "increased instance buffer size to {}", - instance_buffer_pool.buffer_size - ); - } + let command_buffer = match self.render_frame(scene, drawable.texture(), viewport_size) { + Ok(command_buffer) => command_buffer, + Err(error) => { + log::error!("failed to render: {error:#}"); + return; } + }; + + if self.presents_with_transaction { + command_buffer.commit(); + command_buffer.wait_until_scheduled(); + drawable.present(); + } else { + command_buffer.present_drawable(drawable); + command_buffer.commit(); } } + fn render_frame( + &mut self, + scene: &Scene, + texture: &metal::TextureRef, + viewport_size: Size, + ) -> Result { + let mut writer = InstanceBufferWriter::new( + &self.device, + &self.instance_buffer_pool, + self.is_unified_memory, + ); + let instance_bindings = write_instances(scene, &mut writer).with_context(|| { + format!( + "scene too large: {} paths, {} shadows, {} quads, {} underlines, {} mono, {} poly, {} surfaces", + scene.paths.len(), + scene.shadows.len(), + scene.quads.len(), + scene.underlines.len(), + scene.monochrome_sprites.len(), + scene.polychrome_sprites.len(), + scene.surfaces.len(), + ) + })?; + let command_buffer = self.draw_primitives_to_texture( + scene, + &instance_bindings, + &mut writer, + texture, + viewport_size, + )?; + + let instance_buffer_pool = self.instance_buffer_pool.clone(); + let instance_buffer = Cell::new(Some(writer.finish())); + let block = ConcreteBlock::new(move |_| { + if let Some(instance_buffer) = instance_buffer.take() { + instance_buffer_pool.lock().release(instance_buffer); + } + }); + let block = block.copy(); + command_buffer.add_completed_handler(&block); + + Ok(command_buffer) + } + /// Renders the scene to a texture and returns the pixel data as an RGBA image. /// This does not present the frame to screen - useful for visual testing /// where we want to capture what would be rendered without displaying it. @@ -541,83 +552,13 @@ impl MetalRenderer { .next_drawable() .ok_or_else(|| anyhow::anyhow!("Failed to get drawable for render_to_image"))?; - loop { - let mut instance_buffer = self - .instance_buffer_pool - .lock() - .acquire(&self.device, self.is_unified_memory); + let command_buffer = self.render_frame(scene, drawable.texture(), viewport_size)?; - let command_buffer = - self.draw_primitives(scene, &mut instance_buffer, drawable, viewport_size); + // Commit and wait for completion without presenting + command_buffer.commit(); + command_buffer.wait_until_completed(); - match command_buffer { - Ok(command_buffer) => { - let instance_buffer_pool = self.instance_buffer_pool.clone(); - let instance_buffer = Cell::new(Some(instance_buffer)); - let block = ConcreteBlock::new(move |_| { - if let Some(instance_buffer) = instance_buffer.take() { - instance_buffer_pool.lock().release(instance_buffer); - } - }); - let block = block.copy(); - command_buffer.add_completed_handler(&block); - - // Commit and wait for completion without presenting - command_buffer.commit(); - command_buffer.wait_until_completed(); - - // Read pixels from the texture - let texture = drawable.texture(); - let width = texture.width() as u32; - let height = texture.height() as u32; - let bytes_per_row = width as usize * 4; - let buffer_size = height as usize * bytes_per_row; - - let mut pixels = vec![0u8; buffer_size]; - - let region = metal::MTLRegion { - origin: metal::MTLOrigin { x: 0, y: 0, z: 0 }, - size: metal::MTLSize { - width: width as u64, - height: height as u64, - depth: 1, - }, - }; - - texture.get_bytes( - pixels.as_mut_ptr() as *mut std::ffi::c_void, - bytes_per_row as u64, - region, - 0, - ); - - // Convert BGRA to RGBA (swap B and R channels) - for chunk in pixels.chunks_exact_mut(4) { - chunk.swap(0, 2); - } - - return RgbaImage::from_raw(width, height, pixels).ok_or_else(|| { - anyhow::anyhow!("Failed to create RgbaImage from pixel data") - }); - } - Err(err) => { - log::error!( - "failed to render: {}. retrying with larger instance buffer size", - err - ); - let mut instance_buffer_pool = self.instance_buffer_pool.lock(); - let buffer_size = instance_buffer_pool.buffer_size; - if buffer_size >= 256 * 1024 * 1024 { - anyhow::bail!("instance buffer size grew too large: {}", buffer_size); - } - instance_buffer_pool.reset(buffer_size * 2); - log::info!( - "increased instance buffer size to {}", - instance_buffer_pool.buffer_size - ); - } - } - } + read_texture_to_image(drawable.texture()) } /// Renders a scene to an image without requiring a window or CAMetalLayer. @@ -647,92 +588,22 @@ impl MetalRenderer { texture_descriptor.set_storage_mode(metal::MTLStorageMode::Managed); let target_texture = self.device.new_texture(&texture_descriptor); - loop { - let mut instance_buffer = self - .instance_buffer_pool - .lock() - .acquire(&self.device, self.is_unified_memory); + let command_buffer = self.render_frame(scene, &target_texture, size)?; - let command_buffer = - self.draw_primitives_to_texture(scene, &mut instance_buffer, &target_texture, size); - - match command_buffer { - Ok(command_buffer) => { - let instance_buffer_pool = self.instance_buffer_pool.clone(); - let instance_buffer = Cell::new(Some(instance_buffer)); - let block = ConcreteBlock::new(move |_| { - if let Some(instance_buffer) = instance_buffer.take() { - instance_buffer_pool.lock().release(instance_buffer); - } - }); - let block = block.copy(); - command_buffer.add_completed_handler(&block); - - // On discrete GPUs (non-unified memory), Managed textures - // require an explicit blit synchronize before the CPU can - // read back the rendered data. Without this, get_bytes - // returns stale zeros. - if !self.is_unified_memory { - let blit = command_buffer.new_blit_command_encoder(); - blit.synchronize_resource(&target_texture); - blit.end_encoding(); - } - - // Commit and wait for completion - command_buffer.commit(); - command_buffer.wait_until_completed(); - - // Read pixels from the texture - let width = size.width.0 as u32; - let height = size.height.0 as u32; - let bytes_per_row = width as usize * 4; - let buffer_size = height as usize * bytes_per_row; - - let mut pixels = vec![0u8; buffer_size]; - - let region = metal::MTLRegion { - origin: metal::MTLOrigin { x: 0, y: 0, z: 0 }, - size: metal::MTLSize { - width: width as u64, - height: height as u64, - depth: 1, - }, - }; - - target_texture.get_bytes( - pixels.as_mut_ptr() as *mut std::ffi::c_void, - bytes_per_row as u64, - region, - 0, - ); - - // Convert BGRA to RGBA (swap B and R channels) - for chunk in pixels.chunks_exact_mut(4) { - chunk.swap(0, 2); - } - - return RgbaImage::from_raw(width, height, pixels).ok_or_else(|| { - anyhow::anyhow!("Failed to create RgbaImage from pixel data") - }); - } - Err(err) => { - log::error!( - "failed to render: {}. retrying with larger instance buffer size", - err - ); - let mut instance_buffer_pool = self.instance_buffer_pool.lock(); - let buffer_size = instance_buffer_pool.buffer_size; - if buffer_size >= 256 * 1024 * 1024 { - anyhow::bail!("instance buffer size grew too large: {}", buffer_size); - } - instance_buffer_pool.reset(buffer_size * 2); - log::info!( - "increased instance buffer size to {}", - instance_buffer_pool.buffer_size - ); - } - } + // On discrete GPUs (non-unified memory), Managed textures require an + // explicit blit synchronize before the CPU can read back the rendered + // data. Without this, get_bytes returns stale zeros. + if !self.is_unified_memory { + let blit = command_buffer.new_blit_command_encoder(); + blit.synchronize_resource(&target_texture); + blit.end_encoding(); } + + // Commit and wait for completion + command_buffer.commit(); + command_buffer.wait_until_completed(); + + read_texture_to_image(&target_texture) } /// Renders a scene to a reused offscreen texture without reading pixels @@ -769,209 +640,131 @@ impl MetalRenderer { .clone() .expect("just ensured the render target exists"); - loop { - let mut instance_buffer = self - .instance_buffer_pool - .lock() - .acquire(&self.device, self.is_unified_memory); + let command_buffer = self.render_frame(scene, &target_texture, size)?; - let command_buffer = - self.draw_primitives_to_texture(scene, &mut instance_buffer, &target_texture, size); - - match command_buffer { - Ok(command_buffer) => { - let instance_buffer_pool = self.instance_buffer_pool.clone(); - let instance_buffer = Cell::new(Some(instance_buffer)); - let block = ConcreteBlock::new(move |_| { - if let Some(instance_buffer) = instance_buffer.take() { - instance_buffer_pool.lock().release(instance_buffer); - } - }); - let block = block.copy(); - command_buffer.add_completed_handler(&block); - - // Commit without waiting, mirroring presentation to a real - // window where the CPU doesn't block on the GPU. - command_buffer.commit(); - return Ok(()); - } - Err(err) => { - log::error!( - "failed to render: {}. retrying with larger instance buffer size", - err - ); - let mut instance_buffer_pool = self.instance_buffer_pool.lock(); - let buffer_size = instance_buffer_pool.buffer_size; - if buffer_size >= 256 * 1024 * 1024 { - anyhow::bail!("instance buffer size grew too large: {}", buffer_size); - } - instance_buffer_pool.reset(buffer_size * 2); - log::info!( - "increased instance buffer size to {}", - instance_buffer_pool.buffer_size - ); - } - } - } - } - - fn draw_primitives( - &mut self, - scene: &Scene, - instance_buffer: &mut InstanceBuffer, - drawable: &metal::MetalDrawableRef, - viewport_size: Size, - ) -> Result { - self.draw_primitives_to_texture(scene, instance_buffer, drawable.texture(), viewport_size) + // Commit without waiting, mirroring presentation to a real window where + // the CPU doesn't block on the GPU. + command_buffer.commit(); + Ok(()) } fn draw_primitives_to_texture( &mut self, scene: &Scene, - instance_buffer: &mut InstanceBuffer, + instance_bindings: &InstanceBindings, + writer: &mut InstanceBufferWriter, texture: &metal::TextureRef, viewport_size: Size, ) -> Result { let command_queue = self.command_queue.clone(); let command_buffer = command_queue.new_command_buffer(); let alpha = if self.opaque { 1. } else { 0. }; - let mut instance_offset = 0; let mut command_encoder = new_command_encoder_for_texture( command_buffer, texture, viewport_size, - |color_attachment| { - color_attachment.set_load_action(metal::MTLLoadAction::Clear); - color_attachment.set_clear_color(metal::MTLClearColor::new(0., 0., 0., alpha)); - }, + Some(metal::MTLClearColor::new(0., 0., 0., alpha)), ); for batch in scene.batches() { - let ok = match batch { - PrimitiveBatch::Shadows(range) => self.draw_shadows( - &scene.shadows[range], - instance_buffer, - &mut instance_offset, - viewport_size, - command_encoder, - ), - PrimitiveBatch::Quads(range) => self.draw_quads( - &scene.quads[range], - instance_buffer, - &mut instance_offset, - viewport_size, - command_encoder, - ), + match batch { + PrimitiveBatch::Shadows(range) => { + self.draw_shadows(range, instance_bindings, viewport_size, command_encoder) + } + PrimitiveBatch::Quads(range) => { + self.draw_quads(range, instance_bindings, viewport_size, command_encoder) + } PrimitiveBatch::Paths(range) => { let paths = &scene.paths[range]; command_encoder.end_encoding(); let did_draw = self.draw_paths_to_intermediate( paths, - instance_buffer, - &mut instance_offset, + writer, viewport_size, command_buffer, - ); + )?; command_encoder = new_command_encoder_for_texture( command_buffer, texture, viewport_size, - |color_attachment| { - color_attachment.set_load_action(metal::MTLLoadAction::Load); - }, + None, ); if did_draw { - self.draw_paths_from_intermediate( + if let Err(error) = self.draw_paths_from_intermediate( paths, - instance_buffer, - &mut instance_offset, + writer, viewport_size, command_encoder, - ) - } else { - false + ) { + command_encoder.end_encoding(); + return Err(error); + } } } - PrimitiveBatch::Underlines(range) => self.draw_underlines( - &scene.underlines[range], - instance_buffer, - &mut instance_offset, - viewport_size, - command_encoder, - ), + PrimitiveBatch::Underlines(range) => { + self.draw_underlines(range, instance_bindings, viewport_size, command_encoder) + } PrimitiveBatch::MonochromeSprites { texture_id, range } => self .draw_monochrome_sprites( texture_id, - &scene.monochrome_sprites[range], - instance_buffer, - &mut instance_offset, + range, + instance_bindings, viewport_size, command_encoder, ), PrimitiveBatch::PolychromeSprites { texture_id, range } => self .draw_polychrome_sprites( texture_id, - &scene.polychrome_sprites[range], - instance_buffer, - &mut instance_offset, + range, + instance_bindings, viewport_size, command_encoder, ), PrimitiveBatch::Surfaces(range) => self.draw_surfaces( - &scene.surfaces[range], - instance_buffer, - &mut instance_offset, + &scene.surfaces[range.clone()], + range.start, + instance_bindings, viewport_size, command_encoder, ), PrimitiveBatch::SubpixelSprites { .. } => unreachable!(), - }; - if !ok { - command_encoder.end_encoding(); - anyhow::bail!( - "scene too large: {} paths, {} shadows, {} quads, {} underlines, {} mono, {} poly, {} surfaces", - scene.paths.len(), - scene.shadows.len(), - scene.quads.len(), - scene.underlines.len(), - scene.monochrome_sprites.len(), - scene.polychrome_sprites.len(), - scene.surfaces.len(), - ); } } command_encoder.end_encoding(); - if !self.is_unified_memory { - // Sync the instance buffer to the GPU - instance_buffer.metal_buffer.did_modify_range(NSRange { - location: 0, - length: instance_offset as NSUInteger, - }); - } - Ok(command_buffer.to_owned()) } fn draw_paths_to_intermediate( &self, paths: &[Path], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + writer: &mut InstanceBufferWriter, viewport_size: Size, command_buffer: &metal::CommandBufferRef, - ) -> bool { + ) -> Result { if paths.is_empty() { - return true; + return Ok(false); } - let Some(intermediate_texture) = &self.path_intermediate_texture else { - return false; - }; + let intermediate_texture = self + .path_intermediate_texture + .as_ref() + .context("missing path intermediate texture")?; + + let mut vertices = Vec::new(); + for path in paths { + vertices.extend(path.vertices.iter().map(|v| PathRasterizationVertex { + xy_position: v.xy_position, + st_position: v.st_position, + color: path.color, + bounds: path.bounds.intersect(&path.content_mask.bounds), + })); + } + let vertex_instance_bindings = writer.write(&vertices)?; let render_pass_descriptor = metal::RenderPassDescriptor::new(); let color_attachment = render_pass_descriptor @@ -992,27 +785,10 @@ impl MetalRenderer { let command_encoder = command_buffer.new_render_command_encoder(render_pass_descriptor); command_encoder.set_render_pipeline_state(&self.paths_rasterization_pipeline_state); - - align_offset(instance_offset); - let mut vertices = Vec::new(); - for path in paths { - vertices.extend(path.vertices.iter().map(|v| PathRasterizationVertex { - xy_position: v.xy_position, - st_position: v.st_position, - color: path.color, - bounds: path.bounds.intersect(&path.content_mask.bounds), - })); - } - let vertices_bytes_len = mem::size_of_val(vertices.as_slice()); - let next_offset = *instance_offset + vertices_bytes_len; - if next_offset > instance_buffer.size { - command_encoder.end_encoding(); - return false; - } command_encoder.set_vertex_buffer( PathRasterizationInputIndex::Vertices as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&vertex_instance_bindings.buffer), + vertex_instance_bindings.offset as u64, ); command_encoder.set_vertex_bytes( PathRasterizationInputIndex::ViewportSize as u64, @@ -1021,41 +797,29 @@ impl MetalRenderer { ); command_encoder.set_fragment_buffer( PathRasterizationInputIndex::Vertices as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&vertex_instance_bindings.buffer), + vertex_instance_bindings.offset as u64, ); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - unsafe { - ptr::copy_nonoverlapping( - vertices.as_ptr() as *const u8, - buffer_contents, - vertices_bytes_len, - ); - } command_encoder.draw_primitives( metal::MTLPrimitiveType::Triangle, 0, vertices.len() as u64, ); - *instance_offset = next_offset; command_encoder.end_encoding(); - true + Ok(true) } fn draw_shadows( &self, - shadows: &[Shadow], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + shadows: Range, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { if shadows.is_empty() { - return true; + return; } - align_offset(instance_offset); command_encoder.set_render_pipeline_state(&self.shadows_pipeline_state); command_encoder.set_vertex_buffer( @@ -1065,60 +829,39 @@ impl MetalRenderer { ); command_encoder.set_vertex_buffer( ShadowInputIndex::Shadows as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.shadows.buffer), + instance_bindings.shadows.offset as u64, ); command_encoder.set_fragment_buffer( ShadowInputIndex::Shadows as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.shadows.buffer), + instance_bindings.shadows.offset as u64, ); - command_encoder.set_vertex_bytes( ShadowInputIndex::ViewportSize as u64, mem::size_of_val(&viewport_size) as u64, &viewport_size as *const Size as *const _, ); - let shadow_bytes_len = mem::size_of_val(shadows); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - - let next_offset = *instance_offset + shadow_bytes_len; - if next_offset > instance_buffer.size { - return false; - } - - unsafe { - ptr::copy_nonoverlapping( - shadows.as_ptr() as *const u8, - buffer_contents, - shadow_bytes_len, - ); - } - - command_encoder.draw_primitives_instanced( + command_encoder.draw_primitives_instanced_base_instance( metal::MTLPrimitiveType::Triangle, 0, 6, shadows.len() as u64, + shadows.start as u64, ); - *instance_offset = next_offset; - true } fn draw_quads( &self, - quads: &[Quad], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + quads: Range, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { if quads.is_empty() { - return true; + return; } - align_offset(instance_offset); command_encoder.set_render_pipeline_state(&self.quads_pipeline_state); command_encoder.set_vertex_buffer( @@ -1128,59 +871,43 @@ impl MetalRenderer { ); command_encoder.set_vertex_buffer( QuadInputIndex::Quads as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.quads.buffer), + instance_bindings.quads.offset as u64, ); command_encoder.set_fragment_buffer( QuadInputIndex::Quads as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.quads.buffer), + instance_bindings.quads.offset as u64, ); - command_encoder.set_vertex_bytes( QuadInputIndex::ViewportSize as u64, mem::size_of_val(&viewport_size) as u64, &viewport_size as *const Size as *const _, ); - let quad_bytes_len = mem::size_of_val(quads); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - - let next_offset = *instance_offset + quad_bytes_len; - if next_offset > instance_buffer.size { - return false; - } - - unsafe { - ptr::copy_nonoverlapping(quads.as_ptr() as *const u8, buffer_contents, quad_bytes_len); - } - - command_encoder.draw_primitives_instanced( + command_encoder.draw_primitives_instanced_base_instance( metal::MTLPrimitiveType::Triangle, 0, 6, quads.len() as u64, + quads.start as u64, ); - *instance_offset = next_offset; - true } fn draw_paths_from_intermediate( &self, paths: &[Path], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + writer: &mut InstanceBufferWriter, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) -> Result<()> { let Some(first_path) = paths.first() else { - return true; - }; - - let Some(ref intermediate_texture) = self.path_intermediate_texture else { - return false; + return Ok(()); }; + let intermediate_texture = self + .path_intermediate_texture + .as_ref() + .context("missing path intermediate texture")?; command_encoder.set_render_pipeline_state(&self.path_sprites_pipeline_state); command_encoder.set_vertex_buffer( @@ -1222,52 +949,32 @@ impl MetalRenderer { sprites = vec![PathSprite { bounds }]; } - align_offset(instance_offset); - let sprite_bytes_len = mem::size_of_val(sprites.as_slice()); - let next_offset = *instance_offset + sprite_bytes_len; - if next_offset > instance_buffer.size { - return false; - } - + let sprite_instance_bindings = writer.write(&sprites)?; command_encoder.set_vertex_buffer( SpriteInputIndex::Sprites as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&sprite_instance_bindings.buffer), + sprite_instance_bindings.offset as u64, ); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - unsafe { - ptr::copy_nonoverlapping( - sprites.as_ptr() as *const u8, - buffer_contents, - sprite_bytes_len, - ); - } - command_encoder.draw_primitives_instanced( metal::MTLPrimitiveType::Triangle, 0, 6, sprites.len() as u64, ); - *instance_offset = next_offset; - - true + Ok(()) } fn draw_underlines( &self, - underlines: &[Underline], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + underlines: Range, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { if underlines.is_empty() { - return true; + return; } - align_offset(instance_offset); command_encoder.set_render_pipeline_state(&self.underlines_pipeline_state); command_encoder.set_vertex_buffer( @@ -1277,69 +984,39 @@ impl MetalRenderer { ); command_encoder.set_vertex_buffer( UnderlineInputIndex::Underlines as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.underlines.buffer), + instance_bindings.underlines.offset as u64, ); command_encoder.set_fragment_buffer( UnderlineInputIndex::Underlines as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.underlines.buffer), + instance_bindings.underlines.offset as u64, ); - command_encoder.set_vertex_bytes( UnderlineInputIndex::ViewportSize as u64, mem::size_of_val(&viewport_size) as u64, &viewport_size as *const Size as *const _, ); - let underline_bytes_len = mem::size_of_val(underlines); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - - let next_offset = *instance_offset + underline_bytes_len; - if next_offset > instance_buffer.size { - return false; - } - - unsafe { - ptr::copy_nonoverlapping( - underlines.as_ptr() as *const u8, - buffer_contents, - underline_bytes_len, - ); - } - - command_encoder.draw_primitives_instanced( + command_encoder.draw_primitives_instanced_base_instance( metal::MTLPrimitiveType::Triangle, 0, 6, underlines.len() as u64, + underlines.start as u64, ); - *instance_offset = next_offset; - true } fn draw_monochrome_sprites( &self, texture_id: AtlasTextureId, - sprites: &[MonochromeSprite], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + sprites: Range, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { if sprites.is_empty() { - return true; - } - align_offset(instance_offset); - - let sprite_bytes_len = mem::size_of_val(sprites); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - - let next_offset = *instance_offset + sprite_bytes_len; - if next_offset > instance_buffer.size { - return false; + return; } let texture = self.sprite_atlas.metal_texture(texture_id); @@ -1355,8 +1032,8 @@ impl MetalRenderer { ); command_encoder.set_vertex_buffer( SpriteInputIndex::Sprites as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.monochrome_sprites.buffer), + instance_bindings.monochrome_sprites.offset as u64, ); command_encoder.set_vertex_bytes( SpriteInputIndex::ViewportSize as u64, @@ -1370,42 +1047,31 @@ impl MetalRenderer { ); command_encoder.set_fragment_buffer( SpriteInputIndex::Sprites as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.monochrome_sprites.buffer), + instance_bindings.monochrome_sprites.offset as u64, ); command_encoder.set_fragment_texture(SpriteInputIndex::AtlasTexture as u64, Some(&texture)); - unsafe { - ptr::copy_nonoverlapping( - sprites.as_ptr() as *const u8, - buffer_contents, - sprite_bytes_len, - ); - } - - command_encoder.draw_primitives_instanced( + command_encoder.draw_primitives_instanced_base_instance( metal::MTLPrimitiveType::Triangle, 0, 6, sprites.len() as u64, + sprites.start as u64, ); - *instance_offset = next_offset; - true } fn draw_polychrome_sprites( &self, texture_id: AtlasTextureId, - sprites: &[PolychromeSprite], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + sprites: Range, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { if sprites.is_empty() { - return true; + return; } - align_offset(instance_offset); let texture = self.sprite_atlas.metal_texture(texture_id); let texture_size = size( @@ -1420,8 +1086,8 @@ impl MetalRenderer { ); command_encoder.set_vertex_buffer( SpriteInputIndex::Sprites as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.polychrome_sprites.buffer), + instance_bindings.polychrome_sprites.offset as u64, ); command_encoder.set_vertex_bytes( SpriteInputIndex::ViewportSize as u64, @@ -1435,59 +1101,50 @@ impl MetalRenderer { ); command_encoder.set_fragment_buffer( SpriteInputIndex::Sprites as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, + Some(&instance_bindings.polychrome_sprites.buffer), + instance_bindings.polychrome_sprites.offset as u64, ); command_encoder.set_fragment_texture(SpriteInputIndex::AtlasTexture as u64, Some(&texture)); - let sprite_bytes_len = mem::size_of_val(sprites); - let buffer_contents = - unsafe { (instance_buffer.metal_buffer.contents() as *mut u8).add(*instance_offset) }; - - let next_offset = *instance_offset + sprite_bytes_len; - if next_offset > instance_buffer.size { - return false; - } - - unsafe { - ptr::copy_nonoverlapping( - sprites.as_ptr() as *const u8, - buffer_contents, - sprite_bytes_len, - ); - } - - command_encoder.draw_primitives_instanced( + command_encoder.draw_primitives_instanced_base_instance( metal::MTLPrimitiveType::Triangle, 0, 6, sprites.len() as u64, + sprites.start as u64, ); - *instance_offset = next_offset; - true } fn draw_surfaces( &mut self, surfaces: &[PaintSurface], - instance_buffer: &mut InstanceBuffer, - instance_offset: &mut usize, + first_surface: usize, + instance_bindings: &InstanceBindings, viewport_size: Size, command_encoder: &metal::RenderCommandEncoderRef, - ) -> bool { + ) { + if surfaces.is_empty() { + return; + } + command_encoder.set_render_pipeline_state(&self.surfaces_pipeline_state); command_encoder.set_vertex_buffer( SurfaceInputIndex::Vertices as u64, Some(&self.unit_vertices), 0, ); + command_encoder.set_vertex_buffer( + SurfaceInputIndex::Surfaces as u64, + Some(&instance_bindings.surfaces.buffer), + instance_bindings.surfaces.offset as u64, + ); command_encoder.set_vertex_bytes( SurfaceInputIndex::ViewportSize as u64, mem::size_of_val(&viewport_size) as u64, &viewport_size as *const Size as *const _, ); - for surface in surfaces { + for (index, surface) in surfaces.iter().enumerate() { let texture_size = size( DevicePixels::from(surface.image_buffer.get_width() as i32), DevicePixels::from(surface.image_buffer.get_height() as i32), @@ -1521,17 +1178,6 @@ impl MetalRenderer { ) .unwrap(); - align_offset(instance_offset); - let next_offset = *instance_offset + mem::size_of::(); - if next_offset > instance_buffer.size { - return false; - } - - command_encoder.set_vertex_buffer( - SurfaceInputIndex::Surfaces as u64, - Some(&instance_buffer.metal_buffer), - *instance_offset as u64, - ); command_encoder.set_vertex_bytes( SurfaceInputIndex::TextureSize as u64, mem::size_of_val(&texture_size) as u64, @@ -1547,23 +1193,14 @@ impl MetalRenderer { Some(metal::TextureRef::from_ptr(texture as *mut _)) }); - unsafe { - let buffer_contents = (instance_buffer.metal_buffer.contents() as *mut u8) - .add(*instance_offset) - as *mut SurfaceBounds; - ptr::write( - buffer_contents, - SurfaceBounds { - bounds: surface.bounds, - content_mask: surface.content_mask, - }, - ); - } - - command_encoder.draw_primitives(metal::MTLPrimitiveType::Triangle, 0, 6); - *instance_offset = next_offset; + command_encoder.draw_primitives_instanced_base_instance( + metal::MTLPrimitiveType::Triangle, + 0, + 6, + 1, + (first_surface + index) as u64, + ); } - true } } @@ -1571,7 +1208,7 @@ fn new_command_encoder_for_texture<'a>( command_buffer: &'a metal::CommandBufferRef, texture: &'a metal::TextureRef, viewport_size: Size, - configure_color_attachment: impl Fn(&RenderPassColorAttachmentDescriptorRef), + clear_color: Option, ) -> &'a metal::RenderCommandEncoderRef { let render_pass_descriptor = metal::RenderPassDescriptor::new(); let color_attachment = render_pass_descriptor @@ -1580,7 +1217,12 @@ fn new_command_encoder_for_texture<'a>( .unwrap(); color_attachment.set_texture(Some(texture)); color_attachment.set_store_action(metal::MTLStoreAction::Store); - configure_color_attachment(color_attachment); + if let Some(clear_color) = clear_color { + color_attachment.set_load_action(metal::MTLLoadAction::Clear); + color_attachment.set_clear_color(clear_color); + } else { + color_attachment.set_load_action(metal::MTLLoadAction::Load); + } let command_encoder = command_buffer.new_render_command_encoder(render_pass_descriptor); command_encoder.set_viewport(metal::MTLViewport { @@ -1594,6 +1236,36 @@ fn new_command_encoder_for_texture<'a>( command_encoder } +#[cfg(any(test, feature = "test-support"))] +fn read_texture_to_image(texture: &metal::TextureRef) -> Result { + let width = texture.width() as u32; + let height = texture.height() as u32; + let bytes_per_row = width as usize * 4; + let mut pixels = vec![0u8; height as usize * bytes_per_row]; + + let region = metal::MTLRegion { + origin: metal::MTLOrigin { x: 0, y: 0, z: 0 }, + size: metal::MTLSize { + width: width as u64, + height: height as u64, + depth: 1, + }, + }; + texture.get_bytes( + pixels.as_mut_ptr() as *mut std::ffi::c_void, + bytes_per_row as u64, + region, + 0, + ); + + // Convert BGRA to RGBA (swap B and R channels) + for chunk in pixels.chunks_exact_mut(4) { + chunk.swap(0, 2); + } + + RgbaImage::from_raw(width, height, pixels).context("failed to create RgbaImage from pixel data") +} + fn build_pipeline_state( device: &metal::DeviceRef, library: &metal::LibraryRef, @@ -1701,9 +1373,164 @@ fn build_path_rasterization_pipeline_state( .expect("could not create render pipeline state") } -// Align to multiples of 256 make Metal happy. -fn align_offset(offset: &mut usize) { - *offset = (*offset).div_ceil(256) * 256; +#[derive(Clone)] +struct InstanceBinding { + buffer: metal::Buffer, + offset: usize, +} + +struct InstanceBindings { + quads: InstanceBinding, + shadows: InstanceBinding, + underlines: InstanceBinding, + monochrome_sprites: InstanceBinding, + polychrome_sprites: InstanceBinding, + surfaces: InstanceBinding, +} + +fn write_instances(scene: &Scene, writer: &mut InstanceBufferWriter) -> Result { + Ok(InstanceBindings { + quads: writer.write(&scene.quads)?, + shadows: writer.write(&scene.shadows)?, + underlines: writer.write(&scene.underlines)?, + monochrome_sprites: writer.write(&scene.monochrome_sprites)?, + polychrome_sprites: writer.write(&scene.polychrome_sprites)?, + surfaces: writer.write_iter(scene.surfaces.iter().map(|surface| SurfaceBounds { + bounds: surface.bounds, + content_mask: surface.content_mask, + }))?, + }) +} + +struct InstanceBufferWriter { + device: metal::Device, + pool: Arc>, + unified_memory: bool, + filled: Vec<(InstanceBuffer, usize)>, + current: InstanceBuffer, + offset: usize, +} + +impl InstanceBufferWriter { + fn new( + device: &metal::Device, + pool: &Arc>, + unified_memory: bool, + ) -> Self { + let current = pool.lock().acquire(device, unified_memory); + Self { + device: device.clone(), + pool: pool.clone(), + unified_memory, + filled: Vec::new(), + current, + offset: 0, + } + } + + fn allocate(&mut self, count: usize) -> Result<(InstanceBinding, &mut [MaybeUninit])> { + let size = mem::size_of::() * count; + let mut offset = self.offset.next_multiple_of(INSTANCE_BUFFER_ALIGNMENT); + if offset + size > self.current.size { + self.grow(size)?; + offset = 0; + } + self.offset = offset + size; + + let binding = InstanceBinding { + buffer: self.current.metal_buffer.clone(), + offset, + }; + // Safety: the reservation lies within a buffer this frame owns + // exclusively, and never overlaps one handed out earlier. + let values = unsafe { + let start = (self.current.metal_buffer.contents() as *mut u8).add(offset); + slice::from_raw_parts_mut(start.cast::>(), count) + }; + Ok((binding, values)) + } + + fn write(&mut self, values: &[T]) -> Result { + let (binding, destination) = self.allocate::(values.len())?; + unsafe { + ptr::copy_nonoverlapping( + values.as_ptr(), + destination.as_mut_ptr().cast::(), + values.len(), + ); + } + Ok(binding) + } + + fn write_iter( + &mut self, + values: impl ExactSizeIterator, + ) -> Result { + let (binding, destination) = self.allocate::(values.len())?; + for (slot, value) in destination.iter_mut().zip(values) { + slot.write(value); + } + Ok(binding) + } + + fn grow(&mut self, required: usize) -> Result<()> { + let mut pool = self.pool.lock(); + let buffer_size = (pool.buffer_size * 2) + .max(required.next_power_of_two()) + .min(MAX_INSTANCE_BUFFER_SIZE); + anyhow::ensure!( + buffer_size >= required, + "instance buffer needs {required} bytes, above the maximum of {MAX_INSTANCE_BUFFER_SIZE}" + ); + anyhow::ensure!( + buffer_size > self.current.size, + "frame instance data exceeds the {MAX_INSTANCE_BUFFER_SIZE}-byte maximum" + ); + if buffer_size != pool.buffer_size { + log::info!("increased instance buffer size to {buffer_size}"); + pool.reset(buffer_size); + } + let buffer = pool.acquire(&self.device, self.unified_memory); + drop(pool); + + let filled = mem::replace(&mut self.current, buffer); + self.filled.push((filled, self.offset)); + self.offset = 0; + Ok(()) + } + + fn finish(self) -> InstanceBuffer { + let Self { + unified_memory, + filled, + current, + offset, + .. + } = self; + + if !unified_memory { + for (buffer, written) in &filled { + if *written == 0 { + continue; + } + buffer.metal_buffer.did_modify_range(NSRange { + location: 0, + length: *written as NSUInteger, + }); + } + if offset > 0 { + current.metal_buffer.did_modify_range(NSRange { + location: 0, + length: offset as NSUInteger, + }); + } + } + + // Metal retains encoded resources until the command buffer completes. + // Only the final, largest buffer is worth keeping in the pool. + drop(filled); + current + } } #[repr(C)] diff --git a/crates/gpui_wgpu/src/shaders.wgsl b/crates/gpui_wgpu/src/shaders.wgsl index 747a34d81e..c009ed9807 100644 --- a/crates/gpui_wgpu/src/shaders.wgsl +++ b/crates/gpui_wgpu/src/shaders.wgsl @@ -93,8 +93,8 @@ struct GammaParams { @group(0) @binding(0) var globals: GlobalParams; @group(0) @binding(1) var gamma_params: GammaParams; -@group(1) @binding(1) var t_sprite: texture_2d; -@group(1) @binding(2) var s_sprite: sampler; +@group(2) @binding(0) var t_sprite: texture_2d; +@group(2) @binding(1) var s_sprite: sampler; const M_PI_F: f32 = 3.1415926; const GRAYSCALE_FACTORS: vec3 = vec3(0.2126, 0.7152, 0.0722); diff --git a/crates/gpui_wgpu/src/wgpu_renderer.rs b/crates/gpui_wgpu/src/wgpu_renderer.rs index 587315debf..e30ade6e39 100644 --- a/crates/gpui_wgpu/src/wgpu_renderer.rs +++ b/crates/gpui_wgpu/src/wgpu_renderer.rs @@ -1,18 +1,21 @@ use crate::{CompositorGpuHint, WgpuAtlas, WgpuContext}; +use anyhow::{Context as _, Result}; use bytemuck::{Pod, Zeroable}; use gpui::{ - AtlasTextureId, Background, Bounds, DevicePixels, GpuSpecs, MonochromeSprite, Path, Point, - PolychromeSprite, PrimitiveBatch, Quad, ScaledPixels, Scene, Shadow, Size, SubpixelSprite, - Underline, get_gamma_correction_ratios, + AtlasTextureId, Background, Bounds, DevicePixels, GpuSpecs, Path, Point, PrimitiveBatch, + ScaledPixels, Scene, Size, get_gamma_correction_ratios, }; use log::warn; #[cfg(not(target_family = "wasm"))] use raw_window_handle::{HasDisplayHandle, HasWindowHandle}; use std::cell::RefCell; use std::num::NonZeroU64; +use std::ops::Range; use std::rc::Rc; use std::sync::{Arc, Mutex}; +const MAX_INSTANCE_BUFFER_SIZE: u64 = 256 * 1024 * 1024; + #[repr(C)] #[derive(Clone, Copy, Pod, Zeroable)] struct GlobalParams { @@ -94,10 +97,19 @@ struct WgpuPipelines { surfaces: wgpu::RenderPipeline, } +struct InstanceBindings { + quads: wgpu::BindGroup, + shadows: wgpu::BindGroup, + underlines: wgpu::BindGroup, + monochrome_sprites: wgpu::BindGroup, + subpixel_sprites: wgpu::BindGroup, + polychrome_sprites: wgpu::BindGroup, +} + struct WgpuBindGroupLayouts { globals: wgpu::BindGroupLayout, instances: wgpu::BindGroupLayout, - instances_with_texture: wgpu::BindGroupLayout, + texture: wgpu::BindGroupLayout, surfaces: wgpu::BindGroupLayout, } @@ -381,9 +393,15 @@ impl WgpuRenderer { mapped_at_creation: false, }); - let max_buffer_size = device.limits().max_buffer_size; + // Every frame allocation is exposed as one storage-buffer binding, so + // its backing buffer must satisfy both the allocation and binding limits. + let max_buffer_size = device + .limits() + .max_buffer_size + .min(device.limits().max_storage_buffer_binding_size) + .min(MAX_INSTANCE_BUFFER_SIZE); let storage_buffer_alignment = device.limits().min_storage_buffer_offset_alignment as u64; - let initial_instance_buffer_capacity = 2 * 1024 * 1024; + let initial_instance_buffer_capacity = (2 * 1024 * 1024).min(max_buffer_size); let instance_buffer = device.create_buffer(&wgpu::BufferDescriptor { label: Some("instance_buffer"), size: initial_instance_buffer_capacity, @@ -539,29 +557,27 @@ impl WgpuRenderer { entries: &[storage_buffer_entry(0)], }); - let instances_with_texture = - device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - label: Some("instances_with_texture_layout"), - entries: &[ - storage_buffer_entry(0), - wgpu::BindGroupLayoutEntry { - binding: 1, - visibility: wgpu::ShaderStages::VERTEX_FRAGMENT, - ty: wgpu::BindingType::Texture { - sample_type: wgpu::TextureSampleType::Float { filterable: true }, - view_dimension: wgpu::TextureViewDimension::D2, - multisampled: false, - }, - count: None, + let texture = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("texture_layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::VERTEX_FRAGMENT, + ty: wgpu::BindingType::Texture { + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + view_dimension: wgpu::TextureViewDimension::D2, + multisampled: false, }, - wgpu::BindGroupLayoutEntry { - binding: 2, - visibility: wgpu::ShaderStages::VERTEX_FRAGMENT, - ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), - count: None, - }, - ], - }); + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::VERTEX_FRAGMENT, + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + ], + }); let surfaces = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("surfaces_layout"), @@ -610,7 +626,7 @@ impl WgpuRenderer { WgpuBindGroupLayouts { globals, instances, - instances_with_texture, + texture, surfaces, } } @@ -682,13 +698,16 @@ impl WgpuRenderer { fs_entry: &str, globals_layout: &wgpu::BindGroupLayout, data_layout: &wgpu::BindGroupLayout, + texture_layout: Option<&wgpu::BindGroupLayout>, topology: wgpu::PrimitiveTopology, color_targets: &[Option], sample_count: u32, module: &wgpu::ShaderModule| { + let mut bind_group_layouts = vec![Some(globals_layout), Some(data_layout)]; + bind_group_layouts.extend(texture_layout.map(Some)); let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some(&format!("{name}_layout")), - bind_group_layouts: &[Some(globals_layout), Some(data_layout)], + bind_group_layouts: &bind_group_layouts, immediate_size: 0, }); @@ -733,6 +752,7 @@ impl WgpuRenderer { "fs_quad", &layouts.globals, &layouts.instances, + None, wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target.clone())], 1, @@ -745,6 +765,7 @@ impl WgpuRenderer { "fs_shadow", &layouts.globals, &layouts.instances, + None, wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target.clone())], 1, @@ -757,6 +778,7 @@ impl WgpuRenderer { "fs_path_rasterization", &layouts.globals, &layouts.instances, + None, wgpu::PrimitiveTopology::TriangleList, &[Some(wgpu::ColorTargetState { format: surface_format, @@ -785,7 +807,8 @@ impl WgpuRenderer { "vs_path", "fs_path", &layouts.globals, - &layouts.instances_with_texture, + &layouts.instances, + Some(&layouts.texture), wgpu::PrimitiveTopology::TriangleStrip, &[Some(wgpu::ColorTargetState { format: surface_format, @@ -802,6 +825,7 @@ impl WgpuRenderer { "fs_underline", &layouts.globals, &layouts.instances, + None, wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target.clone())], 1, @@ -813,7 +837,8 @@ impl WgpuRenderer { "vs_mono_sprite", "fs_mono_sprite", &layouts.globals, - &layouts.instances_with_texture, + &layouts.instances, + Some(&layouts.texture), wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target.clone())], 1, @@ -839,7 +864,8 @@ impl WgpuRenderer { "vs_subpixel_sprite", "fs_subpixel_sprite", &layouts.globals, - &layouts.instances_with_texture, + &layouts.instances, + Some(&layouts.texture), wgpu::PrimitiveTopology::TriangleStrip, &[Some(wgpu::ColorTargetState { format: surface_format, @@ -858,7 +884,8 @@ impl WgpuRenderer { "vs_poly_sprite", "fs_poly_sprite", &layouts.globals, - &layouts.instances_with_texture, + &layouts.instances, + Some(&layouts.texture), wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target.clone())], 1, @@ -871,6 +898,7 @@ impl WgpuRenderer { "fs_surface", &layouts.globals, &layouts.surfaces, + None, wgpu::PrimitiveTopology::TriangleStrip, &[Some(color_target)], 1, @@ -1202,327 +1230,249 @@ impl WgpuRenderer { ); } - loop { - let mut instance_offset: u64 = 0; - let mut overflow = false; + if let Err(error) = self.record_frame(scene, &frame_view) { + log::error!("{error:#}"); + self.resources().queue.submit(std::iter::empty()); + return false; + } - let mut encoder = - self.resources() - .device - .create_command_encoder(&wgpu::CommandEncoderDescriptor { - label: Some("main_encoder"), - }); + frame.present(); + true + } - { - let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: Some("main_pass"), - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &frame_view, - resolve_target: None, - ops: wgpu::Operations { - load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), - store: wgpu::StoreOp::Store, - }, - depth_slice: None, - })], - depth_stencil_attachment: None, - ..Default::default() + fn record_frame(&mut self, scene: &Scene, frame_view: &wgpu::TextureView) -> Result<()> { + let mut instance_offset = 0; + let instance_bindings = self + .write_instances(scene, &mut instance_offset) + .with_context(|| { + format!( + "scene too large: {} paths, {} shadows, {} quads, {} underlines, {} monochrome sprites, {} subpixel sprites, {} polychrome sprites", + scene.paths.len(), + scene.shadows.len(), + scene.quads.len(), + scene.underlines.len(), + scene.monochrome_sprites.len(), + scene.subpixel_sprites.len(), + scene.polychrome_sprites.len(), + ) + })?; + + let mut encoder = + self.resources() + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("main_encoder"), }); - for batch in scene.batches() { - let ok = match batch { - PrimitiveBatch::Quads(range) => { - self.draw_quads(&scene.quads[range], &mut instance_offset, &mut pass) + { + let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("main_pass"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: frame_view, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT), + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); + + for batch in scene.batches() { + match batch { + PrimitiveBatch::Quads(range) => self.draw_instances( + &instance_bindings.quads, + &self.resources().pipelines.quads, + instance_range(range), + &mut pass, + ), + PrimitiveBatch::Shadows(range) => self.draw_instances( + &instance_bindings.shadows, + &self.resources().pipelines.shadows, + instance_range(range), + &mut pass, + ), + PrimitiveBatch::Paths(range) => { + let paths = &scene.paths[range]; + if paths.is_empty() { + continue; } - PrimitiveBatch::Shadows(range) => self.draw_shadows( - &scene.shadows[range], + + drop(pass); + let rasterized = self.draw_paths_to_intermediate( + &mut encoder, + paths, &mut instance_offset, - &mut pass, - ), - PrimitiveBatch::Paths(range) => { - let paths = &scene.paths[range]; - if paths.is_empty() { - continue; - } + )?; - drop(pass); + pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: Some("main_pass_continued"), + color_attachments: &[Some(wgpu::RenderPassColorAttachment { + view: frame_view, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: wgpu::StoreOp::Store, + }, + depth_slice: None, + })], + depth_stencil_attachment: None, + ..Default::default() + }); - let did_draw = self.draw_paths_to_intermediate( - &mut encoder, + if rasterized { + self.draw_paths_from_intermediate( paths, &mut instance_offset, - ); - - pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { - label: Some("main_pass_continued"), - color_attachments: &[Some(wgpu::RenderPassColorAttachment { - view: &frame_view, - resolve_target: None, - ops: wgpu::Operations { - load: wgpu::LoadOp::Load, - store: wgpu::StoreOp::Store, - }, - depth_slice: None, - })], - depth_stencil_attachment: None, - ..Default::default() - }); - - if did_draw { - self.draw_paths_from_intermediate( - paths, - &mut instance_offset, - &mut pass, - ) - } else { - false - } + &mut pass, + )?; } - PrimitiveBatch::Underlines(range) => self.draw_underlines( - &scene.underlines[range], - &mut instance_offset, - &mut pass, - ), - PrimitiveBatch::MonochromeSprites { texture_id, range } => self - .draw_monochrome_sprites( - &scene.monochrome_sprites[range], - texture_id, - &mut instance_offset, - &mut pass, - ), - PrimitiveBatch::SubpixelSprites { texture_id, range } => self - .draw_subpixel_sprites( - &scene.subpixel_sprites[range], - texture_id, - &mut instance_offset, - &mut pass, - ), - PrimitiveBatch::PolychromeSprites { texture_id, range } => self - .draw_polychrome_sprites( - &scene.polychrome_sprites[range], - texture_id, - &mut instance_offset, - &mut pass, - ), - PrimitiveBatch::Surfaces(_surfaces) => { - // Surfaces are macOS-only for video playback - // Not implemented for Linux/wgpu - true - } - }; - if !ok { - overflow = true; - break; } + PrimitiveBatch::Underlines(range) => self.draw_instances( + &instance_bindings.underlines, + &self.resources().pipelines.underlines, + instance_range(range), + &mut pass, + ), + PrimitiveBatch::MonochromeSprites { texture_id, range } => self.draw_sprites( + &instance_bindings.monochrome_sprites, + texture_id, + &self.resources().pipelines.mono_sprites, + instance_range(range), + &mut pass, + ), + PrimitiveBatch::SubpixelSprites { texture_id, range } => { + let resources = self.resources(); + self.draw_sprites( + &instance_bindings.subpixel_sprites, + texture_id, + resources + .pipelines + .subpixel_sprites + .as_ref() + .unwrap_or(&resources.pipelines.mono_sprites), + instance_range(range), + &mut pass, + ); + } + PrimitiveBatch::PolychromeSprites { texture_id, range } => self.draw_sprites( + &instance_bindings.polychrome_sprites, + texture_id, + &self.resources().pipelines.poly_sprites, + instance_range(range), + &mut pass, + ), + // Surfaces are macOS-only for video playback and are not + // implemented by the WGPU renderer. + PrimitiveBatch::Surfaces(_surfaces) => {} } } - - if overflow { - drop(encoder); - if self.instance_buffer_capacity >= self.max_buffer_size { - log::error!( - "instance buffer size grew too large: {}", - self.instance_buffer_capacity - ); - frame.present(); - return true; - } - self.grow_instance_buffer(); - continue; - } - - self.resources() - .queue - .submit(std::iter::once(encoder.finish())); - frame.present(); - return true; } + + self.resources() + .queue + .submit(std::iter::once(encoder.finish())); + Ok(()) } - fn draw_quads( - &self, - quads: &[Quad], + fn write_instances( + &mut self, + scene: &Scene, instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let data = unsafe { Self::instance_bytes(quads) }; - self.draw_instances( - data, - quads.len() as u32, - &self.resources().pipelines.quads, - instance_offset, - pass, - ) + ) -> Result { + Ok(InstanceBindings { + quads: self.write_instance_binding("quads_bind_group", instance_offset, unsafe { + Self::instance_bytes(&scene.quads) + })?, + shadows: self.write_instance_binding( + "shadows_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&scene.shadows) }, + )?, + underlines: self.write_instance_binding( + "underlines_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&scene.underlines) }, + )?, + monochrome_sprites: self.write_instance_binding( + "monochrome_sprites_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&scene.monochrome_sprites) }, + )?, + subpixel_sprites: self.write_instance_binding( + "subpixel_sprites_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&scene.subpixel_sprites) }, + )?, + polychrome_sprites: self.write_instance_binding( + "polychrome_sprites_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&scene.polychrome_sprites) }, + )?, + }) } - fn draw_shadows( + fn create_texture_bind_group( &self, - shadows: &[Shadow], - instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let data = unsafe { Self::instance_bytes(shadows) }; - self.draw_instances( - data, - shadows.len() as u32, - &self.resources().pipelines.shadows, - instance_offset, - pass, - ) - } - - fn draw_underlines( - &self, - underlines: &[Underline], - instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let data = unsafe { Self::instance_bytes(underlines) }; - self.draw_instances( - data, - underlines.len() as u32, - &self.resources().pipelines.underlines, - instance_offset, - pass, - ) - } - - fn draw_monochrome_sprites( - &self, - sprites: &[MonochromeSprite], - texture_id: AtlasTextureId, - instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let tex_info = self.atlas.get_texture_info(texture_id); - let data = unsafe { Self::instance_bytes(sprites) }; - self.draw_instances_with_texture( - data, - sprites.len() as u32, - &tex_info.view, - &self.resources().pipelines.mono_sprites, - instance_offset, - pass, - ) - } - - fn draw_subpixel_sprites( - &self, - sprites: &[SubpixelSprite], - texture_id: AtlasTextureId, - instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let tex_info = self.atlas.get_texture_info(texture_id); - let data = unsafe { Self::instance_bytes(sprites) }; + label: &str, + texture_view: &wgpu::TextureView, + ) -> wgpu::BindGroup { let resources = self.resources(); - let pipeline = resources - .pipelines - .subpixel_sprites - .as_ref() - .unwrap_or(&resources.pipelines.mono_sprites); - self.draw_instances_with_texture( - data, - sprites.len() as u32, - &tex_info.view, - pipeline, - instance_offset, - pass, - ) - } - - fn draw_polychrome_sprites( - &self, - sprites: &[PolychromeSprite], - texture_id: AtlasTextureId, - instance_offset: &mut u64, - pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - let tex_info = self.atlas.get_texture_info(texture_id); - let data = unsafe { Self::instance_bytes(sprites) }; - self.draw_instances_with_texture( - data, - sprites.len() as u32, - &tex_info.view, - &self.resources().pipelines.poly_sprites, - instance_offset, - pass, - ) + resources + .device + .create_bind_group(&wgpu::BindGroupDescriptor { + label: Some(label), + layout: &resources.bind_group_layouts.texture, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(texture_view), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::Sampler(&resources.atlas_sampler), + }, + ], + }) } fn draw_instances( &self, - data: &[u8], - instance_count: u32, + bind_group: &wgpu::BindGroup, pipeline: &wgpu::RenderPipeline, - instance_offset: &mut u64, + instances: Range, pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - if instance_count == 0 { - return true; + ) { + if instances.is_empty() { + return; } - let Some((offset, size)) = self.write_to_instance_buffer(instance_offset, data) else { - return false; - }; - let resources = self.resources(); - let bind_group = resources - .device - .create_bind_group(&wgpu::BindGroupDescriptor { - label: None, - layout: &resources.bind_group_layouts.instances, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: self.instance_binding(offset, size), - }], - }); pass.set_pipeline(pipeline); - pass.set_bind_group(0, &resources.globals_bind_group, &[]); - pass.set_bind_group(1, &bind_group, &[]); - pass.draw(0..4, 0..instance_count); - true + pass.set_bind_group(0, &self.resources().globals_bind_group, &[]); + pass.set_bind_group(1, bind_group, &[]); + pass.draw(0..4, instances); } - fn draw_instances_with_texture( + fn draw_sprites( &self, - data: &[u8], - instance_count: u32, - texture_view: &wgpu::TextureView, + sprite_instances: &wgpu::BindGroup, + texture_id: AtlasTextureId, pipeline: &wgpu::RenderPipeline, - instance_offset: &mut u64, + instances: Range, pass: &mut wgpu::RenderPass<'_>, - ) -> bool { - if instance_count == 0 { - return true; + ) { + if instances.is_empty() { + return; } - let Some((offset, size)) = self.write_to_instance_buffer(instance_offset, data) else { - return false; - }; - let resources = self.resources(); - let bind_group = resources - .device - .create_bind_group(&wgpu::BindGroupDescriptor { - label: None, - layout: &resources.bind_group_layouts.instances_with_texture, - entries: &[ - wgpu::BindGroupEntry { - binding: 0, - resource: self.instance_binding(offset, size), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: wgpu::BindingResource::TextureView(texture_view), - }, - wgpu::BindGroupEntry { - binding: 2, - resource: wgpu::BindingResource::Sampler(&resources.atlas_sampler), - }, - ], - }); + let texture_info = self.atlas.get_texture_info(texture_id); + let texture = + self.create_texture_bind_group("atlas_texture_bind_group", &texture_info.view); pass.set_pipeline(pipeline); - pass.set_bind_group(0, &resources.globals_bind_group, &[]); - pass.set_bind_group(1, &bind_group, &[]); - pass.draw(0..4, 0..instance_count); - true + pass.set_bind_group(0, &self.resources().globals_bind_group, &[]); + pass.set_bind_group(1, sprite_instances, &[]); + pass.set_bind_group(2, &texture, &[]); + pass.draw(0..4, instances); } unsafe fn instance_bytes(instances: &[T]) -> &[u8] { @@ -1535,11 +1485,11 @@ impl WgpuRenderer { } fn draw_paths_from_intermediate( - &self, + &mut self, paths: &[Path], instance_offset: &mut u64, pass: &mut wgpu::RenderPass<'_>, - ) -> bool { + ) -> Result<()> { let first_path = &paths[0]; let sprites: Vec = if paths.last().map(|p| &p.order) == Some(&first_path.order) { @@ -1557,28 +1507,32 @@ impl WgpuRenderer { vec![PathSprite { bounds }] }; - let resources = self.resources(); - let Some(path_intermediate_view) = resources.path_intermediate_view.as_ref() else { - return true; + let Some(path_intermediate_view) = self.resources().path_intermediate_view.clone() else { + return Ok(()); }; - - let sprite_data = unsafe { Self::instance_bytes(&sprites) }; - self.draw_instances_with_texture( - sprite_data, - sprites.len() as u32, - path_intermediate_view, - &resources.pipelines.paths, - instance_offset, - pass, - ) + let instances = + self.write_instance_binding("path_sprites_bind_group", instance_offset, unsafe { + Self::instance_bytes(&sprites) + })?; + let texture = self.create_texture_bind_group( + "path_intermediate_texture_bind_group", + &path_intermediate_view, + ); + let resources = self.resources(); + pass.set_pipeline(&resources.pipelines.paths); + pass.set_bind_group(0, &resources.globals_bind_group, &[]); + pass.set_bind_group(1, &instances, &[]); + pass.set_bind_group(2, &texture, &[]); + pass.draw(0..4, 0..sprites.len() as u32); + Ok(()) } fn draw_paths_to_intermediate( - &self, + &mut self, encoder: &mut wgpu::CommandEncoder, paths: &[Path], instance_offset: &mut u64, - ) -> bool { + ) -> Result { let mut vertices = Vec::new(); for path in paths { let bounds = path.clipped_bounds(); @@ -1591,30 +1545,18 @@ impl WgpuRenderer { } if vertices.is_empty() { - return true; + return Ok(false); } - let vertex_data = unsafe { Self::instance_bytes(&vertices) }; - let Some((vertex_offset, vertex_size)) = - self.write_to_instance_buffer(instance_offset, vertex_data) - else { - return false; - }; + let data_bind_group = self.write_instance_binding( + "path_rasterization_bind_group", + instance_offset, + unsafe { Self::instance_bytes(&vertices) }, + )?; let resources = self.resources(); - let data_bind_group = resources - .device - .create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("path_rasterization_bind_group"), - layout: &resources.bind_group_layouts.instances, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: self.instance_binding(vertex_offset, vertex_size), - }], - }); - let Some(path_intermediate_view) = resources.path_intermediate_view.as_ref() else { - return true; + return Ok(false); }; let (target_view, resolve_target) = if let Some(ref msaa_view) = resources.path_msaa_view { @@ -1645,46 +1587,75 @@ impl WgpuRenderer { pass.draw(0..vertices.len() as u32, 0..1); } - true + Ok(true) } - fn grow_instance_buffer(&mut self) { - let new_capacity = (self.instance_buffer_capacity * 2).min(self.max_buffer_size); - log::info!("increased instance buffer size to {}", new_capacity); + fn write_instance_binding( + &mut self, + label: &str, + instance_offset: &mut u64, + data: &[u8], + ) -> Result { + // wgpu rejects zero-sized bindings, so empty primitive arrays still + // reserve the 16-byte minimum. + let size = (data.len() as u64).max(16); + let mut offset = (*instance_offset).next_multiple_of(self.storage_buffer_alignment.max(1)); + if offset + size > self.instance_buffer_capacity { + self.grow_instance_buffer(size)?; + offset = 0; + } + *instance_offset = offset + size; + + let resources = self.resources(); + if !data.is_empty() { + resources + .queue + .write_buffer(&resources.instance_buffer, offset, data); + } + let bind_group = resources + .device + .create_bind_group(&wgpu::BindGroupDescriptor { + label: Some(label), + layout: &resources.bind_group_layouts.instances, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { + buffer: &resources.instance_buffer, + offset, + size: NonZeroU64::new(size), + }), + }], + }); + Ok(bind_group) + } + + fn grow_instance_buffer(&mut self, required: u64) -> Result<()> { + let capacity = (self.instance_buffer_capacity * 2) + .max(required.next_power_of_two()) + .min(self.max_buffer_size); + anyhow::ensure!( + capacity >= required, + "instance buffer needs {required} bytes, above the device maximum of {}", + self.max_buffer_size + ); + anyhow::ensure!( + capacity > self.instance_buffer_capacity, + "frame instance data exceeds the {}-byte maximum", + self.max_buffer_size + ); + log::debug!( + "instance buffer grown from {} to {capacity}", + self.instance_buffer_capacity + ); let resources = self.resources_mut(); resources.instance_buffer = resources.device.create_buffer(&wgpu::BufferDescriptor { label: Some("instance_buffer"), - size: new_capacity, + size: capacity, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); - self.instance_buffer_capacity = new_capacity; - } - - fn write_to_instance_buffer( - &self, - instance_offset: &mut u64, - data: &[u8], - ) -> Option<(u64, NonZeroU64)> { - let offset = (*instance_offset).next_multiple_of(self.storage_buffer_alignment); - let size = (data.len() as u64).max(16); - if offset + size > self.instance_buffer_capacity { - return None; - } - let resources = self.resources(); - resources - .queue - .write_buffer(&resources.instance_buffer, offset, data); - *instance_offset = offset + size; - Some((offset, NonZeroU64::new(size).expect("size is at least 16"))) - } - - fn instance_binding(&self, offset: u64, size: NonZeroU64) -> wgpu::BindingResource<'_> { - wgpu::BindingResource::Buffer(wgpu::BufferBinding { - buffer: &self.resources().instance_buffer, - offset, - size: Some(size), - }) + self.instance_buffer_capacity = capacity; + Ok(()) } /// Mark the surface as unconfigured so rendering is skipped until a new @@ -1851,6 +1822,10 @@ impl WgpuRenderer { } } +fn instance_range(range: Range) -> Range { + range.start as u32..range.end as u32 +} + #[cfg(not(target_family = "wasm"))] fn create_surface( instance: &wgpu::Instance, diff --git a/crates/gpui_windows/src/directx_renderer.rs b/crates/gpui_windows/src/directx_renderer.rs index 8a350682e2..5719c12f2d 100644 --- a/crates/gpui_windows/src/directx_renderer.rs +++ b/crates/gpui_windows/src/directx_renderer.rs @@ -27,6 +27,7 @@ pub(crate) const DISABLE_DIRECT_COMPOSITION: &str = "GPUI_DISABLE_DIRECT_COMPOSI const RENDER_TARGET_FORMAT: DXGI_FORMAT = DXGI_FORMAT_B8G8R8A8_UNORM; // This configuration is used for MSAA rendering on paths only, and it's guaranteed to be supported by DirectX 11. const PATH_MULTISAMPLE_COUNT: u32 = 4; +const MAX_INSTANCE_BUFFER_SIZE: usize = 256 * 1024 * 1024; pub(crate) struct FontInfo { pub gamma_ratios: [f32; 4], @@ -95,6 +96,7 @@ struct DirectXRenderPipelines { struct DirectXGlobalElements { global_params_buffer: Option, + batch_params_buffer: Option, sampler: Option, } @@ -229,6 +231,12 @@ impl DirectXRenderer { device_context .OMSetRenderTargets(Some(slice::from_ref(&resources.render_target_view)), None); device_context.RSSetViewports(Some(slice::from_ref(&resources.viewport))); + device_context + .VSSetConstantBuffers(0, Some(slice::from_ref(&self.globals.global_params_buffer))); + device_context + .VSSetConstantBuffers(1, Some(slice::from_ref(&self.globals.batch_params_buffer))); + device_context + .PSSetConstantBuffers(0, Some(slice::from_ref(&self.globals.global_params_buffer))); } Ok(()) } @@ -487,17 +495,11 @@ impl DirectXRenderer { } let devices = self.devices.as_ref().context("devices missing")?; self.pipelines.shadow_pipeline.draw_range( - &devices.device, &devices.device_context, - slice::from_ref( - &self - .resources - .as_ref() - .context("resources missing")? - .viewport, - ), - slice::from_ref(&self.globals.global_params_buffer), - 4, + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, start as u32, len as u32, ) @@ -509,17 +511,11 @@ impl DirectXRenderer { } let devices = self.devices.as_ref().context("devices missing")?; self.pipelines.quad_pipeline.draw_range( - &devices.device, &devices.device_context, - slice::from_ref( - &self - .resources - .as_ref() - .context("resources missing")? - .viewport, - ), - slice::from_ref(&self.globals.global_params_buffer), - 4, + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, start as u32, len as u32, ) @@ -565,8 +561,6 @@ impl DirectXRenderer { self.pipelines.path_rasterization_pipeline.draw( &devices.device_context, - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST, vertices.len() as u32, 1, @@ -629,8 +623,6 @@ impl DirectXRenderer { self.pipelines.path_sprite_pipeline.draw_with_texture( &devices.device_context, slice::from_ref(&resources.path_intermediate_srv), - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), slice::from_ref(&self.globals.sampler), sprites.len() as u32, ) @@ -641,13 +633,12 @@ impl DirectXRenderer { return Ok(()); } let devices = self.devices.as_ref().context("devices missing")?; - let resources = self.resources.as_ref().context("resources missing")?; self.pipelines.underline_pipeline.draw_range( - &devices.device, &devices.device_context, - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), - 4, + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, start as u32, len as u32, ) @@ -663,14 +654,14 @@ impl DirectXRenderer { return Ok(()); } let devices = self.devices.as_ref().context("devices missing")?; - let resources = self.resources.as_ref().context("resources missing")?; let texture_view = self.atlas.get_texture_view(texture_id); self.pipelines.mono_sprites.draw_range_with_texture( - &devices.device, &devices.device_context, &texture_view, - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, slice::from_ref(&self.globals.sampler), start as u32, len as u32, @@ -687,14 +678,14 @@ impl DirectXRenderer { return Ok(()); } let devices = self.devices.as_ref().context("devices missing")?; - let resources = self.resources.as_ref().context("resources missing")?; let texture_view = self.atlas.get_texture_view(texture_id); self.pipelines.subpixel_sprites.draw_range_with_texture( - &devices.device, &devices.device_context, &texture_view, - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, slice::from_ref(&self.globals.sampler), start as u32, len as u32, @@ -711,14 +702,14 @@ impl DirectXRenderer { return Ok(()); } let devices = self.devices.as_ref().context("devices missing")?; - let resources = self.resources.as_ref().context("resources missing")?; let texture_view = self.atlas.get_texture_view(texture_id); self.pipelines.poly_sprites.draw_range_with_texture( - &devices.device, &devices.device_context, &texture_view, - slice::from_ref(&resources.viewport), - slice::from_ref(&self.globals.global_params_buffer), + self.globals + .batch_params_buffer + .as_ref() + .context("batch params buffer missing")?, slice::from_ref(&self.globals.sampler), start as u32, len as u32, @@ -948,18 +939,8 @@ impl DirectComposition { impl DirectXGlobalElements { pub fn new(device: &ID3D11Device) -> Result { - let global_params_buffer = unsafe { - let desc = D3D11_BUFFER_DESC { - ByteWidth: std::mem::size_of::() as u32, - Usage: D3D11_USAGE_DYNAMIC, - BindFlags: D3D11_BIND_CONSTANT_BUFFER.0 as u32, - CPUAccessFlags: D3D11_CPU_ACCESS_WRITE.0 as u32, - ..Default::default() - }; - let mut buffer = None; - device.CreateBuffer(&desc, None, Some(&mut buffer))?; - buffer - }; + let global_params_buffer = create_constant_buffer::(device)?; + let batch_params_buffer = create_constant_buffer::(device)?; let sampler = unsafe { let desc = D3D11_SAMPLER_DESC { @@ -981,6 +962,7 @@ impl DirectXGlobalElements { Ok(Self { global_params_buffer, + batch_params_buffer, sampler, }) } @@ -997,6 +979,15 @@ struct GlobalParams { _pad: [u32; 3], } +#[derive(Clone, Copy, Debug, Default)] +#[repr(C, align(16))] +struct BatchParams { + start_index: u32, + _padding: [u32; 3], +} + +const _: () = assert!(std::mem::size_of::() == 16); + struct PipelineState { label: &'static str, vertex: ID3D11VertexShader, @@ -1046,7 +1037,17 @@ impl PipelineState { data: &[T], ) -> Result<()> { if self.buffer_size < data.len() { - let new_buffer_size = data.len().next_power_of_two(); + let element_size = std::mem::size_of::(); + let required_size = std::mem::size_of_val(data); + anyhow::ensure!( + required_size <= MAX_INSTANCE_BUFFER_SIZE, + "{} buffer needs {required_size} bytes, above the maximum of {MAX_INSTANCE_BUFFER_SIZE}", + self.label + ); + let new_buffer_size = data + .len() + .next_power_of_two() + .min(MAX_INSTANCE_BUFFER_SIZE / element_size); log::debug!( "Updating {} buffer size from {} to {}", self.label, @@ -1065,8 +1066,6 @@ impl PipelineState { fn draw( &self, device_context: &ID3D11DeviceContext, - viewport: &[D3D11_VIEWPORT], - global_params: &[Option], topology: D3D_PRIMITIVE_TOPOLOGY, vertex_count: u32, instance_count: u32, @@ -1075,10 +1074,8 @@ impl PipelineState { device_context, slice::from_ref(&self.view), topology, - viewport, &self.vertex, &self.fragment, - global_params, &self.blend_state, ); unsafe { @@ -1091,8 +1088,6 @@ impl PipelineState { &self, device_context: &ID3D11DeviceContext, texture: &[Option], - viewport: &[D3D11_VIEWPORT], - global_params: &[Option], sampler: &[Option], instance_count: u32, ) -> Result<()> { @@ -1100,10 +1095,8 @@ impl PipelineState { device_context, slice::from_ref(&self.view), D3D_PRIMITIVE_TOPOLOGY_TRIANGLESTRIP, - viewport, &self.vertex, &self.fragment, - global_params, &self.blend_state, ); unsafe { @@ -1118,51 +1111,52 @@ impl PipelineState { fn draw_range( &self, - device: &ID3D11Device, device_context: &ID3D11DeviceContext, - viewport: &[D3D11_VIEWPORT], - global_params: &[Option], - vertex_count: u32, + batch_params_buffer: &ID3D11Buffer, first_instance: u32, instance_count: u32, ) -> Result<()> { - let view = create_buffer_view_range(device, &self.buffer, first_instance, instance_count)?; + anyhow::ensure!( + first_instance as usize + instance_count as usize <= self.buffer_size, + "DirectX instance range exceeds the {} buffer", + self.label + ); + update_batch_start(device_context, batch_params_buffer, first_instance)?; set_pipeline_state( device_context, - slice::from_ref(&view), + slice::from_ref(&self.view), D3D_PRIMITIVE_TOPOLOGY_TRIANGLESTRIP, - viewport, &self.vertex, &self.fragment, - global_params, &self.blend_state, ); unsafe { - device_context.DrawInstanced(vertex_count, instance_count, 0, 0); + device_context.DrawInstanced(4, instance_count, 0, 0); } Ok(()) } fn draw_range_with_texture( &self, - device: &ID3D11Device, device_context: &ID3D11DeviceContext, texture: &[Option], - viewport: &[D3D11_VIEWPORT], - global_params: &[Option], + batch_params_buffer: &ID3D11Buffer, sampler: &[Option], first_instance: u32, instance_count: u32, ) -> Result<()> { - let view = create_buffer_view_range(device, &self.buffer, first_instance, instance_count)?; + anyhow::ensure!( + first_instance as usize + instance_count as usize <= self.buffer_size, + "DirectX instance range exceeds the {} buffer", + self.label + ); + update_batch_start(device_context, batch_params_buffer, first_instance)?; set_pipeline_state( device_context, - slice::from_ref(&view), + slice::from_ref(&self.view), D3D_PRIMITIVE_TOPOLOGY_TRIANGLESTRIP, - viewport, &self.vertex, &self.fragment, - global_params, &self.blend_state, ); unsafe { @@ -1282,7 +1276,14 @@ fn create_resources( create_path_intermediate_texture(&devices.device, width, height)?; let (path_intermediate_msaa_texture, path_intermediate_msaa_view) = create_path_intermediate_msaa_texture_and_view(&devices.device, width, height)?; - let viewport = set_viewport(&devices.device_context, width as f32, height as f32); + let viewport = D3D11_VIEWPORT { + TopLeftX: 0.0, + TopLeftY: 0.0, + Width: width as f32, + Height: height as f32, + MinDepth: 0.0, + MaxDepth: 1.0, + }; Ok(( render_target, render_target_view, @@ -1369,20 +1370,6 @@ fn create_path_intermediate_msaa_texture_and_view( Ok((msaa_texture, Some(msaa_view.unwrap()))) } -#[inline] -fn set_viewport(device_context: &ID3D11DeviceContext, width: f32, height: f32) -> D3D11_VIEWPORT { - let viewport = [D3D11_VIEWPORT { - TopLeftX: 0.0, - TopLeftY: 0.0, - Width: width, - Height: height, - MinDepth: 0.0, - MaxDepth: 1.0, - }]; - unsafe { device_context.RSSetViewports(Some(&viewport)) }; - viewport[0] -} - #[inline] fn set_rasterizer_state(device: &ID3D11Device, device_context: &ID3D11DeviceContext) -> Result<()> { let desc = D3D11_RASTERIZER_DESC { @@ -1504,6 +1491,22 @@ fn create_fragment_shader(device: &ID3D11Device, bytes: &[u8]) -> Result(device: &ID3D11Device) -> Result> { + const { assert!(std::mem::size_of::() != 0 && std::mem::size_of::().is_multiple_of(16)) }; + let desc = D3D11_BUFFER_DESC { + ByteWidth: std::mem::size_of::() as u32, + Usage: D3D11_USAGE_DYNAMIC, + BindFlags: D3D11_BIND_CONSTANT_BUFFER.0 as u32, + CPUAccessFlags: D3D11_CPU_ACCESS_WRITE.0 as u32, + MiscFlags: 0, + StructureByteStride: 0, + }; + let mut buffer = None; + unsafe { device.CreateBuffer(&desc, None, Some(&mut buffer)) }?; + Ok(buffer) +} + #[inline] fn create_buffer( device: &ID3D11Device, @@ -1533,32 +1536,6 @@ fn create_buffer_view( Ok(view) } -#[inline] -fn create_buffer_view_range( - device: &ID3D11Device, - buffer: &ID3D11Buffer, - first_element: u32, - num_elements: u32, -) -> Result> { - let desc = D3D11_SHADER_RESOURCE_VIEW_DESC { - Format: DXGI_FORMAT_UNKNOWN, - ViewDimension: D3D11_SRV_DIMENSION_BUFFER, - Anonymous: D3D11_SHADER_RESOURCE_VIEW_DESC_0 { - Buffer: D3D11_BUFFER_SRV { - Anonymous1: D3D11_BUFFER_SRV_0 { - FirstElement: first_element, - }, - Anonymous2: D3D11_BUFFER_SRV_1 { - NumElements: num_elements, - }, - }, - }, - }; - let mut view = None; - unsafe { device.CreateShaderResourceView(buffer, Some(&desc), Some(&mut view)) }?; - Ok(view) -} - #[inline] fn update_buffer( device_context: &ID3D11DeviceContext, @@ -1574,26 +1551,37 @@ fn update_buffer( Ok(()) } +#[inline] +fn update_batch_start( + device_context: &ID3D11DeviceContext, + buffer: &ID3D11Buffer, + first_instance: u32, +) -> Result<()> { + update_buffer( + device_context, + buffer, + &[BatchParams { + start_index: first_instance, + _padding: [0; 3], + }], + ) +} + #[inline] fn set_pipeline_state( device_context: &ID3D11DeviceContext, buffer_view: &[Option], topology: D3D_PRIMITIVE_TOPOLOGY, - viewport: &[D3D11_VIEWPORT], vertex_shader: &ID3D11VertexShader, fragment_shader: &ID3D11PixelShader, - global_params: &[Option], blend_state: &ID3D11BlendState, ) { unsafe { device_context.VSSetShaderResources(1, Some(buffer_view)); device_context.PSSetShaderResources(1, Some(buffer_view)); device_context.IASetPrimitiveTopology(topology); - device_context.RSSetViewports(Some(viewport)); device_context.VSSetShader(vertex_shader, None); device_context.PSSetShader(fragment_shader, None); - device_context.VSSetConstantBuffers(0, Some(global_params)); - device_context.PSSetConstantBuffers(0, Some(global_params)); device_context.OMSetBlendState(blend_state, None, 0xFFFFFFFF); } } diff --git a/crates/gpui_windows/src/shaders.hlsl b/crates/gpui_windows/src/shaders.hlsl index 483f039938..d20cb317db 100644 --- a/crates/gpui_windows/src/shaders.hlsl +++ b/crates/gpui_windows/src/shaders.hlsl @@ -9,6 +9,11 @@ cbuffer GlobalParams: register(b0) { uint3 global_pad; }; +cbuffer BatchParams: register(b1) { + uint batch_start_index; + uint3 batch_pad; +}; + Texture2D t_sprite: register(t0); SamplerState s_sprite: register(s0); @@ -525,8 +530,9 @@ struct QuadFragmentInput { StructuredBuffer quads: register(t1); -QuadVertexOutput quad_vertex(uint vertex_id: SV_VertexID, uint quad_id: SV_InstanceID) { +QuadVertexOutput quad_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + uint quad_id = batch_start_index + instance_id; Quad quad = quads[quad_id]; float4 device_position = to_device_position(unit_vertex, quad.bounds); @@ -875,8 +881,9 @@ struct ShadowFragmentInput { StructuredBuffer shadows: register(t1); -ShadowVertexOutput shadow_vertex(uint vertex_id: SV_VertexID, uint shadow_id: SV_InstanceID) { +ShadowVertexOutput shadow_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + uint shadow_id = batch_start_index + instance_id; Shadow shadow = shadows[shadow_id]; Bounds bounds; @@ -1080,8 +1087,9 @@ struct UnderlineFragmentInput { StructuredBuffer underlines: register(t1); -UnderlineVertexOutput underline_vertex(uint vertex_id: SV_VertexID, uint underline_id: SV_InstanceID) { +UnderlineVertexOutput underline_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + uint underline_id = batch_start_index + instance_id; Underline underline = underlines[underline_id]; float4 device_position = to_device_position(unit_vertex, underline.bounds); float4 clip_distance = distance_from_clip_rect(unit_vertex, underline.bounds, @@ -1155,8 +1163,9 @@ struct MonochromeSpriteFragmentInput { StructuredBuffer mono_sprites: register(t1); -MonochromeSpriteVertexOutput monochrome_sprite_vertex(uint vertex_id: SV_VertexID, uint sprite_id: SV_InstanceID) { +MonochromeSpriteVertexOutput monochrome_sprite_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + uint sprite_id = batch_start_index + instance_id; MonochromeSprite sprite = mono_sprites[sprite_id]; float4 device_position = to_device_position_transformed(unit_vertex, sprite.bounds, sprite.transformation); @@ -1178,8 +1187,8 @@ float4 monochrome_sprite_fragment(MonochromeSpriteFragmentInput input): SV_Targe return float4(input.color.rgb, input.color.a * alpha_corrected); } -MonochromeSpriteVertexOutput subpixel_sprite_vertex(uint vertex_id: SV_VertexID, uint sprite_id: SV_InstanceID) { - return monochrome_sprite_vertex(vertex_id, sprite_id); +MonochromeSpriteVertexOutput subpixel_sprite_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { + return monochrome_sprite_vertex(vertex_id, instance_id); } SubpixelSpriteFragmentOutput subpixel_sprite_fragment(MonochromeSpriteFragmentInput input) { @@ -1227,8 +1236,9 @@ struct PolychromeSpriteFragmentInput { StructuredBuffer poly_sprites: register(t1); -PolychromeSpriteVertexOutput polychrome_sprite_vertex(uint vertex_id: SV_VertexID, uint sprite_id: SV_InstanceID) { +PolychromeSpriteVertexOutput polychrome_sprite_vertex(uint vertex_id: SV_VertexID, uint instance_id: SV_InstanceID) { float2 unit_vertex = float2(float(vertex_id & 1u), 0.5 * float(vertex_id & 2u)); + uint sprite_id = batch_start_index + instance_id; PolychromeSprite sprite = poly_sprites[sprite_id]; float4 device_position = to_device_position(unit_vertex, sprite.bounds); float4 clip_distance = distance_from_clip_rect(unit_vertex, sprite.bounds,