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,